diff --git a/src/Activation.cpp b/src/Activation.cpp index 5ef62d5..3c43e68 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -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); diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index c14fa44..e268cb2 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -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; ilayers_n; i++) + for(int i=0; ilayers_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(buf); return r; } - +/* + if(name.find("Route") == 0) { + RouteRT *r = new RouteRT(); + r->in = readBUF(buf); + for(int i=0; ic_in[i] = readBUF(buf); + r->c = readBUF(buf); + r->h = readBUF(buf); + r->w = readBUF(buf); + return r; + } +*/ FatalError("Cant deserialize Plugin"); return NULL; } diff --git a/src/pluginsRT/RouteRT.cpp b/src/pluginsRT/RouteRT.cpp new file mode 100644 index 0000000..b79dac8 --- /dev/null +++ b/src/pluginsRT/RouteRT.cpp @@ -0,0 +1,82 @@ +#include +#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(outputs[0]); + + int offset = 0; + for(int i=0; i(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(buffer); + tk::dnn::writeBUF(buf, in); + for(int i=0; i(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; } diff --git a/src/pluginsRT/UpsampleRT.cpp b/src/pluginsRT/UpsampleRT.cpp index c7e1e25..1272cc9 100644 --- a/src/pluginsRT/UpsampleRT.cpp +++ b/src/pluginsRT/UpsampleRT.cpp @@ -43,8 +43,8 @@ public: dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); dnnType *dstData = reinterpret_cast(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; } diff --git a/src/pluginsRT/YoloRT.cpp b/src/pluginsRT/YoloRT.cpp index ada4f56..548d80a 100644 --- a/src/pluginsRT/YoloRT.cpp +++ b/src/pluginsRT/YoloRT.cpp @@ -57,7 +57,7 @@ public: } } - std::cout<<"YOLO END\n"; + //std::cout<<"YOLO END\n"; return 0; } diff --git a/tests/yolo3_berkeley/yolo3_berkeley.cpp b/tests/yolo3_berkeley/yolo3_berkeley.cpp index b08836d..998c564 100644 --- a/tests/yolo3_berkeley/yolo3_berkeley.cpp +++ b/tests/yolo3_berkeley/yolo3_berkeley.cpp @@ -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; }