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

yolo3 ok

Browse Source
This commit is contained in:
Francesco Gatti
2019-01-04 22:28:10 +01:00
parent 2e8d0b1002
commit 88097a3774
7 changed files with 139 additions and 23 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/Activation.cpp
+1 -1
View File
@@ -47,7 +47,7 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) {
if(act_mode == ACTIVATION_LEAKY) {
activationLEAKYForward(srcData, dstData, dim.tot());
} else {
dnnType alpha = dnnType(1);
dnnType beta = dnnType(0);
src/NetworkRT.cpp
+19 -3
View File
@@ -14,6 +14,7 @@ using namespace nvinfer1;
#include "pluginsRT/ActivationLeakyRT.cpp"
#include "pluginsRT/ReorgRT.cpp"
#include "pluginsRT/RegionRT.cpp"
//#include "pluginsRT/RouteRT.cpp"
#include "pluginsRT/ShortcutRT.cpp"
#include "pluginsRT/YoloRT.cpp"
#include "pluginsRT/UpsampleRT.cpp"
@@ -77,9 +78,10 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
Ilay->setName( (l->getLayerName() + std::to_string(i)).c_str() );
input = Ilay->getOutput(0);
input->setName( (l->getLayerName() + std::to_string(i) + "_out").c_str() );
if(l->getLayerType() == LAYER_YOLO)
networkRT->markOutput(*input);
tensors[l] = input;
}
if(input == NULL)
@@ -308,9 +310,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) {
//std::cout<<"convert route\n";
ITensor **tens = new ITensor*[l->layers_n];
for(int i=0; i<l->layers_n; i++)
for(int i=0; i<l->layers_n; i++) {
tens[i] = tensors[l->layers[i]];
}
IConcatenationLayer *lRT = networkRT->addConcatenation(tens, l->layers_n);
//IPlugin *plugin = new RouteRT();
//IPluginLayer *lRT = networkRT->addPlugin(tens, l->layers_n, *plugin);
checkNULL(lRT);
return lRT;
@@ -445,7 +450,18 @@ public:
r->w = readBUF<int>(buf);
return r;
}
/*
if(name.find("Route") == 0) {
RouteRT *r = new RouteRT();
r->in = readBUF<int>(buf);
for(int i=0; i<RouteRT::MAX_INPUTS; i++)
r->c_in[i] = readBUF<int>(buf);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
return r;
}
*/
FatalError("Cant deserialize Plugin");
return NULL;
}
src/pluginsRT/RouteRT.cpp
+82
View File
@@ -0,0 +1,82 @@
#include<cassert>
#include "kernels.h"
class RouteRT : public IPlugin {
public:
RouteRT() {
}
~RouteRT(){
}
int getNbOutputs() const override {
return 1;
}
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]};
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
in = nbInputs;
c = 0;
for(int i=0; i<nbInputs; i++) {
c_in[i] = inputDims[i].d[0];
c += inputDims[i].d[0];
}
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
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 {
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;
}
return 0;
}
virtual size_t getSerializationSize() override {
return (4+MAX_INPUTS)*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, in);
for(int i=0; i<MAX_INPUTS; i++)
tk::dnn::writeBUF(buf, c_in[i]);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
}
static const int MAX_INPUTS = 4;
int in;
int c_in[MAX_INPUTS];
int c, h, w;
};
src/pluginsRT/ShortcutRT.cpp
+1 -1
View File
@@ -44,7 +44,7 @@ public:
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, c, h, w, 1);
shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, c, h, w, 1, stream);
return 0;
}
src/pluginsRT/UpsampleRT.cpp
+2 -2
View File
@@ -43,8 +43,8 @@ public:
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
fill(dstData, batchSize*c*h*w, 0.0);
upsampleForward(srcData, dstData, batchSize, c, h, w, stride, 1, 1);
fill(dstData, batchSize*c*h*w*stride*stride, 0.0, stream);
upsampleForward(srcData, dstData, batchSize, c, h, w, stride, 1, 1, stream);
return 0;
}
src/pluginsRT/YoloRT.cpp
+1 -1
View File
@@ -57,7 +57,7 @@ public:
}
}
std::cout<<"YOLO END\n";
//std::cout<<"YOLO END\n";
return 0;
}
tests/yolo3_berkeley/yolo3_berkeley.cpp
+33 -15
View File
@@ -77,7 +77,11 @@ const char *c102_bin = "../tests/yolo3_berkeley/layers/c102.bin";
const char *c103_bin = "../tests/yolo3_berkeley/layers/c103.bin";
const char *c104_bin = "../tests/yolo3_berkeley/layers/c104.bin";
const char *c105_bin = "../tests/yolo3_berkeley/layers/c105.bin";
const char *output_bin = "../tests/yolo3_berkeley/debug/layer93_out.bin";
const char *output_bins[3] = {
"../tests/yolo3_berkeley/debug/layer82_out.bin",
"../tests/yolo3_berkeley/debug/layer94_out.bin",
"../tests/yolo3_berkeley/debug/layer106_out.bin"
};
int main() {
@@ -234,7 +238,7 @@ int main() {
tk::dnn::Conv2d c80 (&net,1024, 3, 3, 1, 1, 1, 1, c80_bin, true);
tk::dnn::Activation a80 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c81 (&net, 45, 1, 1, 1, 1, 0, 0, c81_bin, false);
tk::dnn::Yolo y82 (&net, 10, 3);
tk::dnn::Yolo yolo0 (&net, 10, 3);
tk::dnn::Layer *m83_layers[1] = { &a79 };
tk::dnn::Route m83 (&net, m83_layers, 1);
@@ -243,7 +247,7 @@ int main() {
tk::dnn::Upsample u85 (&net, 2);
tk::dnn::Layer *m86_layers[2] = { &u85, &s61 };
tk::dnn::Route m86 (&net, m86_layers, 1); // ROUTE ERROR IN RT INFERENCE
tk::dnn::Route m86 (&net, m86_layers, 2);
tk::dnn::Conv2d c87 (&net, 256, 1, 1, 1, 1, 0, 0, c87_bin, true);
tk::dnn::Activation a87 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c88 (&net, 512, 3, 3, 1, 1, 1, 1, c88_bin, true);
@@ -254,10 +258,11 @@ int main() {
tk::dnn::Activation a90 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c91 (&net, 256, 1, 1, 1, 1, 0, 0, c91_bin, true);
tk::dnn::Activation a91 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c92 (&net, 512, 3, 3, 1, 1, 1, 1, c92_bin, true);
tk::dnn::Activation a92 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c93 (&net, 45, 1, 1, 1, 1, 0, 0, c93_bin, false);
tk::dnn::Yolo y94 (&net, 10, 3);
tk::dnn::Yolo yolo1 (&net, 10, 3);
tk::dnn::Layer *m95_layers[1] = { &a91 };
tk::dnn::Route m95 (&net, m95_layers, 1);
@@ -281,7 +286,7 @@ int main() {
tk::dnn::Conv2d c104 (&net, 256, 3, 3, 1, 1, 1, 1, c104_bin, true);
tk::dnn::Activation a104 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c105 (&net, 45, 1, 1, 1, 1, 0, 0, c105_bin, false);
tk::dnn::Yolo y106 (&net, 10, 3);
tk::dnn::Yolo yolo2 (&net, 10, 3);
// Load input
dnnType *data;
@@ -294,32 +299,45 @@ int main() {
//convert network to tensorRT
tk::dnn::NetworkRT netRT(&net, "yolo3_berkeley.rt");
dnnType *out_data, *out_data2; // cudnn output, tensorRT output
// the network have 3 outputs
tk::dnn::dataDim_t out_dim[3];
out_dim[0] = yolo0.output_dim;
out_dim[1] = yolo1.output_dim;
out_dim[2] = yolo2.output_dim;
dnnType *cudnn_out[3], *rt_out[3];
tk::dnn::dataDim_t dim1 = dim; //input dim
printCenteredTitle(" CUDNN inference ", '=', 30); {
dim1.print();
TIMER_START
out_data = net.infer(dim1, data);
net.infer(dim1, data);
TIMER_STOP
dim1.print();
}
cudnn_out[0] = yolo0.dstData;
cudnn_out[1] = yolo1.dstData;
cudnn_out[2] = yolo2.dstData;
tk::dnn::dataDim_t dim2 = dim;
printCenteredTitle(" TENSORRT inference ", '=', 30); {
dim2.print();
TIMER_START
out_data2 = netRT.infer(dim2, data);
netRT.infer(dim2, data);
TIMER_STOP
dim2.print();
}
rt_out[0] = (dnnType*)netRT.buffersRT[1];
rt_out[1] = (dnnType*)netRT.buffersRT[2];
rt_out[2] = (dnnType*)netRT.buffersRT[3];
printCenteredTitle(" CHECK RESULTS ", '=', 30);
dnnType *out, *out_h;
int out_dim = net.getOutputDim().tot();
//readBinaryFile(output_bin, out_dim, &out_h, &out);
//std::cout<<"CUDNN vs correct"; checkResult(out_dim, out_data, out);
//std::cout<<"TRT vs correct"; checkResult(out_dim, out_data2, out);
std::cout<<"CUDNN vs TRT "; checkResult(out_dim, out_data, out_data2);
for(int i=0; i<3; i++) {
printCenteredTitle((std::string(" YOLO ") + std::to_string(i) + " CHECK RESULTS ").c_str(), '=', 30);
dnnType *out, *out_h;
int odim = out_dim[i].tot();
readBinaryFile(output_bins[i], odim, &out_h, &out);
std::cout<<"CUDNN vs correct"; checkResult(odim, cudnn_out[i], out);
std::cout<<"TRT vs correct"; checkResult(odim, rt_out[i], out);
std::cout<<"CUDNN vs TRT "; checkResult(odim, cudnn_out[i], rt_out[i]);
}
return 0;
}