diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 1947309..0d6d472 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -321,7 +321,7 @@ public: int winH, winW; int strideH, strideW; int paddingH, paddingW; - bool test; + bool maxpoolfixedsize; tkdnnPoolingMode_t pool_mode; Pooling(Network *net, int winH, int winW, @@ -474,7 +474,7 @@ public: dnnType *predictions; - static const int MAX_DETECTIONS = 4096; + static const int MAX_DETECTIONS = 1024; static Yolo::detection *allocateDetections(int nboxes, int classes); static void mergeDetections(Yolo::detection *dets, int ndets, int classes); }; diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 2434696..27c108d 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -36,6 +36,7 @@ using namespace nvinfer1; #include "pluginsRT/DeformableConvRT.h" #include "pluginsRT/FlattenConcatRT.h" #include "pluginsRT/ReshapeRT.h" +#include "pluginsRT/MaxPoolingFixedSizeRT.h" class PluginFactory : IPluginFactory { diff --git a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h new file mode 100644 index 0000000..992042b --- /dev/null +++ b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h @@ -0,0 +1,74 @@ +#include +#include "../kernels.h" + +class MaxPoolFixedSizeRT : public IPlugin { + +public: + MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding) { + this->c = c; + this->h = h; + this->w = w; + this->n = n; + this->stride_H = strideH; + this->stride_W = strideW; + this->winSize = winSize; + this->padding = padding; + } + + ~MaxPoolFixedSizeRT(){ + } + + int getNbOutputs() const override { + return 1; + } + + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return DimsCHW{this->c, this->h, this->w}; + } + + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + } + + 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 { + + std::cout<n<<" "<c<<" "<h<<" "<w<<" "<stride_H<<" "<stride_W<<" "<winSize<<" "<padding<(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + MaxPoolingForward(srcData, dstData, this->n, this->c, this->h, this->w, this->stride_H, this->stride_W, this->winSize, this->padding); + return 0; + } + + + virtual size_t getSerializationSize() override { + return 8*sizeof(int); + } + + virtual void serialize(void* buffer) override { + char *buf = reinterpret_cast(buffer); + + tk::dnn::writeBUF(buf, this->c); + tk::dnn::writeBUF(buf, this->h); + tk::dnn::writeBUF(buf, this->w); + tk::dnn::writeBUF(buf, this->n); + tk::dnn::writeBUF(buf, this->stride_H); + tk::dnn::writeBUF(buf, this->stride_W); + tk::dnn::writeBUF(buf, this->winSize); + tk::dnn::writeBUF(buf, this->padding); + } + + int n, c, h, w; + int stride_H, stride_W; + int winSize; + int padding; +}; diff --git a/include/tkDNN/pluginsRT/ShortcutRT.h b/include/tkDNN/pluginsRT/ShortcutRT.h index ff7b528..fd2c60d 100644 --- a/include/tkDNN/pluginsRT/ShortcutRT.h +++ b/include/tkDNN/pluginsRT/ShortcutRT.h @@ -4,7 +4,10 @@ class ShortcutRT : public IPlugin { public: - ShortcutRT() { + ShortcutRT(tk::dnn::dataDim_t bdim) { + this->bc = bdim.c; + this->bh = bdim.h; + this->bw = bdim.w; } ~ShortcutRT(){ @@ -44,22 +47,27 @@ public: dnnType *dstData = reinterpret_cast(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, stream); + shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, bc, bh, bw, 1, stream); return 0; } virtual size_t getSerializationSize() override { - return 3*sizeof(int); + return 6*sizeof(int); } virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer); + tk::dnn::writeBUF(buf, bc); + tk::dnn::writeBUF(buf, bh); + tk::dnn::writeBUF(buf, bw); tk::dnn::writeBUF(buf, c); tk::dnn::writeBUF(buf, h); tk::dnn::writeBUF(buf, w); + } int c, h, w; + int bc, bh, bw; }; diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 44185f9..711f079 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -304,23 +304,36 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) { if(l->pool_mode == tkdnnPoolingMode_t::POOLING_AVERAGE) ptype = PoolingType::kAVERAGE; if(l->pool_mode == tkdnnPoolingMode_t::POOLING_AVERAGE_EXCLUDE_PADDING) ptype = PoolingType::kMAX_AVERAGE_BLEND; - - if(l->paddingH == 0 && l->paddingW == 0 && l->input_dim.h == l->output_dim.h && l->input_dim.w == l->output_dim.w) + if(l->maxpoolfixedsize) { - IPlugin *plugin = new ResizeLayerRT( l->output_dim.c,l->output_dim.h+1,l->output_dim.w+1 ); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); - checkNULL(lRT); - lRT->setName( "Resize" ); - - input = lRT->getOutput(0); + IPlugin *plugin = new MaxPoolFixedSizeRT(l->output_dim.c, l->output_dim.h, l->output_dim.w, l->output_dim.n, l->strideH, l->strideW, l->winH, l->winH-1); + + IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + + checkNULL(lRT); + lRT->setName( "MaxPoolingFixedSize" ); + return lRT; } + else + { + if(l->paddingH == 0 && l->paddingW == 0 && l->input_dim.h == l->output_dim.h && l->input_dim.w == l->output_dim.w) + { + IPlugin *plugin = new ResizeLayerRT( l->output_dim.c,l->output_dim.h+1,l->output_dim.w+1 ); + IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + checkNULL(lRT); + lRT->setName( "Resize" ); - IPoolingLayer *lRT = networkRT->addPooling(*input, ptype, DimsHW{l->winH, l->winW}); - checkNULL(lRT); + input = lRT->getOutput(0); + } - lRT->setPadding(DimsHW{l->paddingH, l->paddingW}); - lRT->setStride(DimsHW{l->strideH, l->strideW}); - return lRT; + IPoolingLayer *lRT = networkRT->addPooling(*input, ptype, DimsHW{l->winH, l->winW}); + checkNULL(lRT); + + lRT->setPadding(DimsHW{l->paddingH, l->paddingW}); + lRT->setStride(DimsHW{l->strideH, l->strideW}); + return lRT; + + } } ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { @@ -437,18 +450,18 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) { //std::cout<<"New plugin Shortcut\n"; ITensor *back_tens = tensors[l->backLayer]; - /* + // plugin version - IPlugin *plugin = new ShortcutRT(); + IPlugin *plugin = new ShortcutRT(l->backLayer->output_dim); ITensor **inputs = new ITensor*[2]; inputs[0] = input; inputs[1] = back_tens; IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin); checkNULL(lRT); - */ + - IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM); - checkNULL(lRT); + // IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM); + // checkNULL(lRT); return lRT; } @@ -602,13 +615,31 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa } if(name.find("Shortcut") == 0) { - ShortcutRT *r = new ShortcutRT(); + tk::dnn::dataDim_t bdim; + bdim.c = readBUF(buf); + bdim.h = readBUF(buf); + bdim.w = readBUF(buf); + bdim.l = 1; + + ShortcutRT *r = new ShortcutRT(bdim); r->c = readBUF(buf); r->h = readBUF(buf); r->w = readBUF(buf); return r; } + if(name.find("Pooling") == 0) { + MaxPoolFixedSizeRT *r = new MaxPoolFixedSizeRT( readBUF(buf), //c + readBUF(buf), //h + readBUF(buf), //w + readBUF(buf), //n + readBUF(buf), //strideH + readBUF(buf), //strideW + readBUF(buf), //winSize + readBUF(buf)); //padding + return r; + } + if(name.find("Resize") == 0) { ResizeLayerRT *r = new ResizeLayerRT(readBUF(buf), //o_c readBUF(buf), //o_h diff --git a/src/Pooling.cpp b/src/Pooling.cpp index 272edca..48f12cf 100644 --- a/src/Pooling.cpp +++ b/src/Pooling.cpp @@ -7,7 +7,7 @@ namespace tk { namespace dnn { Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW, int paddingH, int paddingW, - tkdnnPoolingMode_t pool_mode, bool final, bool test) : + tkdnnPoolingMode_t pool_mode, bool final, bool maxpoolfixedsize) : Layer(net, final) { this->winH = winH; @@ -17,7 +17,7 @@ Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW, this->pool_mode = pool_mode; this->paddingH = paddingH; this->paddingW = paddingW; - this->test = test; + this->maxpoolfixedsize = maxpoolfixedsize; checkCUDNN( cudnnCreatePoolingDescriptor(&poolingDesc) ); @@ -114,7 +114,7 @@ dnnType* Pooling::infer(dataDim_t &dim, dnnType* srcData) { - if(this->test) + if(this->maxpoolfixedsize) { MaxPoolingForward(poolSrc, poolDst, dim.n, dim.c, dim.h, dim.w, this->strideH, this->strideW, this->winH, this->winH-1); } diff --git a/tests/csresnext50-panet-spp/csresnext50-panet-spp.cfg b/tests/csresnext50-panet-spp/csresnext50-panet-spp.cfg new file mode 100644 index 0000000..ece1122 --- /dev/null +++ b/tests/csresnext50-panet-spp/csresnext50-panet-spp.cfg @@ -0,0 +1,1018 @@ +[net] +# Testing +#batch=1 +#subdivisions=1 +# Training +batch=64 +subdivisions=16 +width=416 +height=416 +channels=3 +momentum=0.9 +decay=0.0005 +angle=0 +saturation = 1.5 +exposure = 1.5 +hue=.1 + +learning_rate=0.001 +burn_in=1000 +max_batches = 500500 +policy=steps +steps=400000,450000 +scales=.1,.1 + +#19:104x104 38:52x52 65:26x26 80:13x13 for 416 + +[convolutional] +batch_normalize=1 +filters=64 +size=7 +stride=2 +pad=1 +activation=leaky + +[maxpool] +size=2 +stride=2 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=leaky + +# 1-1 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 1-2 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 1-3 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 1-T + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -1,-16 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +groups=32 +stride=2 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=linear + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=linear + +# 2-1 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 2-2 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 2-3 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 2-T + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -1,-16 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +groups=32 +stride=2 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=linear + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=linear + +# 3-1 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 3-2 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 3-3 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 3-4 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 3-5 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 3-T + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -1,-24 + +[convolutional] +batch_normalize=1 +filters=1024 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=1024 +size=3 +groups=32 +stride=2 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=1024 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=1024 +size=1 +stride=1 +pad=1 +activation=leaky + +# 4-1 + +[convolutional] +batch_normalize=1 +filters=1024 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=1024 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=1024 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 4-2 + +[convolutional] +batch_normalize=1 +filters=1024 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=1024 +size=3 +groups=32 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=1024 +size=1 +stride=1 +pad=1 +activation=linear + +[shortcut] +from=-4 +activation=leaky + +# 4-T + +[convolutional] +batch_normalize=1 +filters=1024 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -1,-12 + +[convolutional] +batch_normalize=1 +filters=2048 +size=1 +stride=1 +pad=1 +activation=leaky + +########################## + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=1024 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +### SPP ### +[maxpool] +stride=1 +size=5 + +[route] +layers=-2 + +[maxpool] +stride=1 +size=9 + +[route] +layers=-4 + +[maxpool] +stride=1 +size=13 + +[route] +layers=-1,-3,-5,-6 +### End SPP ### + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=1024 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[upsample] +stride=2 + +[route] +layers = 65 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -1, -3 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[upsample] +stride=2 + +[route] +layers = 38 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -1, -3 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +########################## + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=leaky + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=linear + + +[yolo] +mask = 0,1,2 +anchors = 10,13, 16,30, 33,23, 30,61, 62,45, 59,119, 116,90, 156,198, 373,326 +classes=80 +num=9 +jitter=.3 +ignore_thresh = .7 +truth_thresh = 1 +random=1 + +[route] +layers = -4 + +[convolutional] +batch_normalize=1 +size=3 +stride=2 +pad=1 +filters=256 +activation=leaky + +[route] +layers = -1, -16 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=leaky + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=linear + + +[yolo] +mask = 3,4,5 +anchors = 10,13, 16,30, 33,23, 30,61, 62,45, 59,119, 116,90, 156,198, 373,326 +classes=80 +num=9 +jitter=.3 +ignore_thresh = .7 +truth_thresh = 1 +random=1 + +[route] +layers = -4 + +[convolutional] +batch_normalize=1 +size=3 +stride=2 +pad=1 +filters=512 +activation=leaky + +[route] +layers = -1, -37 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=1024 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=1024 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=1024 +activation=leaky + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=linear + + +[yolo] +mask = 6,7,8 +anchors = 10,13, 16,30, 33,23, 30,61, 62,45, 59,119, 116,90, 156,198, 373,326 +classes=80 +num=9 +jitter=.3 +ignore_thresh = .7 +truth_thresh = 1 +random=1 diff --git a/tests/csresnext50-panet-spp/csresnext50-panet-spp.cpp b/tests/csresnext50-panet-spp/csresnext50-panet-spp.cpp index dd4e2fb..087fd71 100644 --- a/tests/csresnext50-panet-spp/csresnext50-panet-spp.cpp +++ b/tests/csresnext50-panet-spp/csresnext50-panet-spp.cpp @@ -339,7 +339,7 @@ int main() tk::dnn::Conv2d c83(&net, 512, 1, 1, 1, 1, 0, 0, c83_bin, true); tk::dnn::Activation a83(&net, tk::dnn::ACTIVATION_LEAKY); - // //SPP + //SPP tk::dnn::Pooling p84(&net, 5, 5, 1, 1,0,0, tk::dnn::POOLING_MAX, false, true); tk::dnn::Layer *r85_layers[1] = {&a83}; tk::dnn::Route r85(&net, r85_layers, 1); @@ -351,7 +351,7 @@ int main() tk::dnn::Pooling p88(&net, 13, 13, 1, 1, 12, 12, tk::dnn::POOLING_MAX, false, true); tk::dnn::Layer *r89_layers[4] = {&p88, &p86, &p84, &a83}; tk::dnn::Route r89(&net, r89_layers, 4); - // //END SPP + //END SPP tk::dnn::Conv2d c90(&net, 512, 1, 1, 1, 1, 0, 0, c90_bin, true); tk::dnn::Activation a90(&net, tk::dnn::ACTIVATION_LEAKY); @@ -454,17 +454,15 @@ int main() tk::dnn::Conv2d c136(&net, 255, 1, 1, 1, 1, 0, 0, c136_bin, false); tk::dnn::Yolo yolo137(&net, classes, 3, g137_bin); - + yolo[0] = &yolo115; + yolo[1] = &yolo126; + yolo[2] = &yolo137; - // yolo[0] = &yolo115; - // yolo[1] = &yolo126; - // yolo[2] = &yolo137; - - // // fill classes names - // for (int i = 0; i < 3; i++) - // { - // yolo[i]->classesNames = {"person", "bicycle", "car", "motorbike", "aeroplane", "bus", "train", "truck", "boat", "traffic light", "fire hydrant", "stop sign", "parking meter", "bench", "bird", "cat", "dog", "horse", "sheep", "cow", "elephant", "bear", "zebra", "giraffe", "backpack", "umbrella", "handbag", "tie", "suitcase", "frisbee", "skis", "snowboard", "sports ball", "kite", "baseball bat", "baseball glove", "skateboard", "surfboard", "tennis racket", "bottle", "wine glass", "cup", "fork", "knife", "spoon", "bowl", "banana", "apple", "sandwich", "orange", "broccoli", "carrot", "hot dog", "pizza", "donut", "cake", "chair", "sofa", "pottedplant", "bed", "diningtable", "toilet", "tvmonitor", "laptop", "mouse", "remote", "keyboard", "cell phone", "microwave", "oven", "toaster", "sink", "refrigerator", "book", "clock", "vase", "scissors", "teddy bear", "hair drier", "toothbrush"}; - // } + // fill classes names + for (int i = 0; i < 3; i++) + { + yolo[i]->classesNames = {"person", "bicycle", "car", "motorbike", "aeroplane", "bus", "train", "truck", "boat", "traffic light", "fire hydrant", "stop sign", "parking meter", "bench", "bird", "cat", "dog", "horse", "sheep", "cow", "elephant", "bear", "zebra", "giraffe", "backpack", "umbrella", "handbag", "tie", "suitcase", "frisbee", "skis", "snowboard", "sports ball", "kite", "baseball bat", "baseball glove", "skateboard", "surfboard", "tennis racket", "bottle", "wine glass", "cup", "fork", "knife", "spoon", "bowl", "banana", "apple", "sandwich", "orange", "broccoli", "carrot", "hot dog", "pizza", "donut", "cake", "chair", "sofa", "pottedplant", "bed", "diningtable", "toilet", "tvmonitor", "laptop", "mouse", "remote", "keyboard", "cell phone", "microwave", "oven", "toaster", "sink", "refrigerator", "book", "clock", "vase", "scissors", "teddy bear", "hair drier", "toothbrush"}; + } // Load input dnnType *data; @@ -475,13 +473,13 @@ int main() net.print(); // //convert network to tensorRT - // tk::dnn::NetworkRT netRT(&net, "csresnext50-panet-spp.rt"); + tk::dnn::NetworkRT netRT(&net, "csresnext50-panet-spp.rt"); // the network have 3 outputs - // tk::dnn::dataDim_t out_dim[3]; - // for (int i = 0; i < 3; i++) - // out_dim[i] = yolo[i]->output_dim; - // dnnType *cudnn_out[3], *rt_out[3]; + tk::dnn::dataDim_t out_dim[3]; + for (int i = 0; i < 3; i++) + out_dim[i] = yolo[i]->output_dim; + dnnType *cudnn_out[3], *rt_out[3]; tk::dnn::dataDim_t dim1 = dim; //input dim printCenteredTitle(" CUDNN inference ", '=', 30); @@ -492,69 +490,62 @@ int main() TIMER_STOP dim1.print(); } - dnnType *cudnn_out = net.layers[net.num_layers-1]->dstData; - tk::dnn::dataDim_t out_dim = net.layers[net.num_layers-1]->output_dim; - dnnType *out1, *out1_h; - int odim1 = out_dim.tot(); - readBinaryFile(output_bin, odim1, &out1_h, &out1); - std::cout << "CUDNN vs correct" << std::endl; - // printDeviceVector(odim1, cudnn_out); - checkResult(odim1, cudnn_out, out1); + + for (int i = 0; i < 3; i++) + cudnn_out[i] = yolo[i]->dstData; - // for (int i = 0; i < 3; i++) - // cudnn_out[i] = yolo[i]->dstData; + printCenteredTitle(" compute detections ", '=', 30); + TIMER_START + int ndets = 0; + tk::dnn::Yolo::detection *dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes); + for (int i = 0; i < 3; i++) + yolo[i]->computeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5); + tk::dnn::Yolo::mergeDetections(dets, ndets, classes); - // printCenteredTitle(" compute detections ", '=', 30); - // TIMER_START - // int ndets = 0; - // tk::dnn::Yolo::detection *dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes); - // for (int i = 0; i < 3; i++) - // yolo[i]->computeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5); - // tk::dnn::Yolo::mergeDetections(dets, ndets, classes); + 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.); - // 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 cl = 0; + for (int c = 0; c < classes; ++c) + { + float prob = dets[j].prob[c]; + if (prob > 0) + cl = c; + } + std::cout << cl << ": " << x0 << " " << y0 << " " << x1 << " " << y1 << "\n"; + } + TIMER_STOP - // int cl = 0; - // for (int c = 0; c < classes; ++c) - // { - // float prob = dets[j].prob[c]; - // if (prob > 0) - // cl = c; - // } - // std::cout << cl << ": " << x0 << " " << y0 << " " << x1 << " " << y1 << "\n"; - // } - // TIMER_STOP + tk::dnn::dataDim_t dim2 = dim; + printCenteredTitle(" TENSORRT inference ", '=', 30); + { + dim2.print(); + TIMER_START + netRT.infer(dim2, data); + TIMER_STOP + dim2.print(); + } - // tk::dnn::dataDim_t dim2 = dim; - // printCenteredTitle(" TENSORRT inference ", '=', 30); - // { - // dim2.print(); - // TIMER_START - // netRT.infer(dim2, data); - // TIMER_STOP - // dim2.print(); - // } - // for (int i = 0; i < 3; i++) - // rt_out[i] = (dnnType *)netRT.buffersRT[i + 1]; + for (int i = 0; i < 3; i++) + rt_out[i] = (dnnType *)netRT.buffersRT[i + 1]; - // 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]); - // } + 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; }