(tkDNN): Make Resnet-101-AP-GeM work with TRT

Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>

CMakeLists.txt

@@ -178,6 +178,9 @@ target_link_libraries(test_mobilenetv2ssd512 tkDNN)
 add_executable(test_resnet101 tests/backbones/resnet101/resnet101.cpp)
 target_link_libraries(test_resnet101 tkDNN)
 
+add_executable(test_Resnet-101-AP-GeM tests/backbones/resnet101/Resnet-101-AP-GeM.cpp)
+target_link_libraries(test_Resnet-101-AP-GeM tkDNN)
+
 add_executable(test_dla34 tests/backbones/dla34/dla34.cpp)
 target_link_libraries(test_dla34 tkDNN)
 

include/tkDNN/Layer.h

@@ -488,7 +488,8 @@ typedef enum {
     POOLING_MAX     = 0,
     POOLING_AVERAGE = 1,                    // count for average includes padded values
     POOLING_AVERAGE_EXCLUDE_PADDING = 2,    // count for average does not include padded values
-    POOLING_MAX_FIXEDSIZE = 100             // max pool darknet fashion
+    POOLING_MAX_FIXEDSIZE = 100,             // max pool darknet fashion
+    POOLING_GENERALIZED_MEAN_P = 200       // mean pooling with pow parameter 
 } tkdnnPoolingMode_t;
 
 /**
@@ -498,6 +499,7 @@ typedef enum {
 class Pooling : public Layer {
 
 public:
+    float pow_param;
     int winH, winW;
     int strideH, strideW;
     int paddingH, paddingW;
@@ -508,7 +510,7 @@ public:
     Pooling(Network *net, int winH, int winW, 
             int strideH, int strideW, 
             int paddingH, int paddingW,
-            tkdnnPoolingMode_t pool_mode);
+            tkdnnPoolingMode_t pool_mode, float p = 1.0f);
     virtual ~Pooling();
     virtual layerType_t getLayerType() { return LAYER_POOLING; };
 

include/tkDNN/kernels.h

@@ -20,6 +20,8 @@ void reorgForward(dnnType *srcData, dnnType *dstData,
 
 void MaxPoolingForward(dnnType *srcData, dnnType *dstData, int n, int c, int h, int w, int stride_x, int stride_y, int size, int padding, cudaStream_t stream = cudaStream_t(0));
 
+void GeneralizedMeanPoolingP(dnnType* srcData, dnnType* dstData, int n, int c, int h, int w, float p, cudaStream_t stream  = cudaStream_t(0));
+
 void softmaxForward(float *input, int n, int batch, int batch_offset,
                     int groups, int group_offset, int stride, float temp, float *output, cudaStream_t stream = cudaStream_t(0));
 

include/tkDNN/pluginsRT/GeneralizedMeanPoolingPRT.h

@@ -0,0 +1,103 @@
+#include<cassert>
+#include "../kernels.h"
+#include <NvInfer.h>
+#include <vector>
+#include <utils.h>
+
+
+namespace nvinfer1 {
+    class GeneralizedMeanPoolingPRT : public IPluginV2Ext {
+
+    public:
+        GeneralizedMeanPoolingPRT(int input_c, int input_h, int input_w, int input_n, int output_c, int output_h, int output_w, int output_n, float p) ;
+
+        GeneralizedMeanPoolingPRT(const void *data, size_t length) ;
+
+        ~GeneralizedMeanPoolingPRT() ;
+
+        int getNbOutputs() const NOEXCEPT override ;
+
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
+
+        int initialize() NOEXCEPT override ;
+
+        void terminate() NOEXCEPT override ;
+
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                            cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR <= 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+
+        size_t getSerializationSize() const NOEXCEPT override ;
+
+        void serialize(void *buffer) const NOEXCEPT override ;
+
+        void destroy() NOEXCEPT override ;
+
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        const char *getPluginType() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        IPluginV2Ext *clone() const NOEXCEPT override ;
+
+        DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+        void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+        bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+        bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+        void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+                              int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+                              bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+                              int32_t maxBatchSize) NOEXCEPT override;
+
+        void detachFromContext() NOEXCEPT override;
+
+
+        int i_n, i_c, i_h, i_w, o_n, o_c, o_h, o_w;
+        float p;
+
+    private:
+        std::string mPluginNamespace;
+    };
+
+    class GeneralizedMeanPoolingPRTPluginCreator : public IPluginCreator {
+    public:
+        GeneralizedMeanPoolingPRTPluginCreator() ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+        IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+        const char *getPluginName() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+    private:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+
+    };
+
+    REGISTER_TENSORRT_PLUGIN(GeneralizedMeanPoolingPRTPluginCreator);
+};

src/NetworkRT.cpp

@@ -514,6 +514,28 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) {
         checkNULL(lRT);
         return lRT;
 
+    }
+    else if(l->pool_mode == tkdnnPoolingMode_t::POOLING_GENERALIZED_MEAN_P)
+    {
+        auto creator = getPluginRegistry()->getPluginCreator("GeneralizedMeanPoolingPRT_tkDNN","1");
+        std::vector<PluginField> mPluginAttributes;
+        PluginFieldCollection mFC{};
+        mPluginAttributes.emplace_back(PluginField("i_c",&l->input_dim.c,PluginFieldType::kINT32,1));
+        mPluginAttributes.emplace_back(PluginField("i_h",&l->input_dim.h,PluginFieldType::kINT32,1));
+        mPluginAttributes.emplace_back(PluginField("i_w",&l->input_dim.w,PluginFieldType::kINT32,1));
+        mPluginAttributes.emplace_back(PluginField("i_n",&l->input_dim.n,PluginFieldType::kINT32,1));
+        mPluginAttributes.emplace_back(PluginField("o_c",&l->output_dim.c,PluginFieldType::kINT32,1));
+        mPluginAttributes.emplace_back(PluginField("o_h",&l->output_dim.h,PluginFieldType::kINT32,1));
+        mPluginAttributes.emplace_back(PluginField("o_w",&l->output_dim.w,PluginFieldType::kINT32,1));
+        mPluginAttributes.emplace_back(PluginField("o_n",&l->output_dim.n,PluginFieldType::kINT32,1));
+        mPluginAttributes.emplace_back(PluginField("p",&l->pow_param,PluginFieldType::kFLOAT32,1));
+        mFC.nbFields = mPluginAttributes.size();
+        mFC.fields = mPluginAttributes.data();
+        auto *plugin = creator->createPlugin(l->getLayerName().c_str(),&mFC);
+        auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
+        checkNULL(lRT);
+        return lRT;
+
     }
     else
     {

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) : 
+                  tkdnnPoolingMode_t pool_mode, float p) : 
     Layer(net) {
 
     this->winH = winH;
@@ -18,6 +18,7 @@ Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW,
     this->paddingH = paddingH;
     this->paddingW = paddingW;
     this->padding = winH -1;
+    this->pow_param = p;
 
     checkCUDNN( cudnnCreatePoolingDescriptor(&poolingDesc) );
 
@@ -42,6 +43,8 @@ Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW,
 
     cudnnPoolingMode_t cudnn_pool_mode = cudnnPoolingMode_t(pool_mode);
     if(pool_mode == POOLING_MAX_FIXEDSIZE) cudnn_pool_mode = cudnnPoolingMode_t(tkdnnPoolingMode_t::POOLING_MAX);
+    if(pool_mode == POOLING_GENERALIZED_MEAN_P) cudnn_pool_mode = cudnnPoolingMode_t(tkdnnPoolingMode_t::POOLING_AVERAGE);
+    
 
     checkCUDNN( cudnnSetPooling2dDescriptor(poolingDesc, cudnn_pool_mode,
                 CUDNN_NOT_PROPAGATE_NAN, winH, winW, paddingH, paddingW, strideH, strideW) );
@@ -114,6 +117,9 @@ dnnType* Pooling::infer(dataDim_t &dim, dnnType* srcData) {
     if(pool_mode == tkdnnPoolingMode_t::POOLING_MAX_FIXEDSIZE){
         MaxPoolingForward(poolSrc, poolDst, dim.n, dim.c, dim.h, dim.w, this->strideH, this->strideW, this->winH, this->winH-1);
     }
+    else if(pool_mode == tkdnnPoolingMode_t::POOLING_GENERALIZED_MEAN_P){
+        GeneralizedMeanPoolingP(poolSrc, poolDst, dim.n, dim.c, dim.h, dim.w, pow_param);
+    }
     else{
         dnnType alpha = dnnType(1);
         dnnType beta = dnnType(0);

src/kernels/pooling.cu

@@ -39,6 +39,43 @@ __global__ void forward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c
     output[out_index] = max;
 }
 
+__global__ void forward_gen_avgpool_p_layer_kernel(int n, int w, int h, int c, float p, float *input, float *output)
+{
+    int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+    if(id >= n) return;
+
+    int k = id % c;
+    id /= c;
+    int b = id;
+
+    int i;
+    int out_index = (k + c*b);
+    output[out_index] = 0;
+    for(i = 0; i < w*h; ++i){
+        int in_index = i + h*w*(k + b*c);
+        float in = 1e-6f;
+        if (input[in_index] > 1e-6f)
+            in = input[in_index];
+        in = pow(in, p);
+        
+        output[out_index] += in;
+    }
+    
+    output[out_index] /= w*h;
+    output[out_index] = pow(output[out_index], 1. / p);
+}
+
+void GeneralizedMeanPoolingP(dnnType* srcData, dnnType* dstData, int n, int c, int h, int w, float p, cudaStream_t stream)
+{
+
+    int tot_size = n*c*h*w;
+
+    int blocks = (tot_size+255)/256;
+    int threads = 256;
+
+    forward_gen_avgpool_p_layer_kernel<<<blocks, threads, 0, stream>>>(n*c, h, w, c, p, srcData, dstData);
+}
+
 void MaxPoolingForward(dnnType* srcData, dnnType* dstData, int n, int c, int h, int w, int stride_x, int stride_y, int size, int padding, cudaStream_t stream) 
 {
 

src/pluginsRT/GeneralizedMeanPoolingPRT.cpp

@@ -0,0 +1,216 @@
+#include <tkDNN/pluginsRT/GeneralizedMeanPoolingPRT.h>
+using namespace nvinfer1;
+
+std::vector<PluginField> GeneralizedMeanPoolingPRTPluginCreator::mPluginAttributes;
+PluginFieldCollection GeneralizedMeanPoolingPRTPluginCreator::mFC{};
+
+static const char* GENERALIZEDMEANPOOLINGPRT_PLUGIN_VERSION{"1"};
+static const char* GENERALIZEDMEANPOOLINGPRT_PLUGIN_NAME{"GeneralizedMeanPoolingPRT_tkDNN"};
+
+GeneralizedMeanPoolingPRT::GeneralizedMeanPoolingPRT(int input_c, int input_h, int input_w, int input_n, int output_c, int output_h, int output_w, int output_n, float p){
+    this->i_c = input_c;
+    this->i_h = input_h;
+    this->i_w = input_w;
+    this->i_n = input_n;
+    this->o_c = output_c;
+    this->o_h = output_h;
+    this->o_w = output_w;
+    this->o_n = output_n;
+    this->p = p;
+}
+
+GeneralizedMeanPoolingPRT::GeneralizedMeanPoolingPRT(const void *data, size_t length) {
+    const char *buf = reinterpret_cast<const char*>(data),*bufCheck = buf;
+    i_c = readBUF<int>(buf);
+    i_h = readBUF<int>(buf);
+    i_w = readBUF<int>(buf);
+    i_n = readBUF<int>(buf);
+    o_c = readBUF<int>(buf);
+    o_h = readBUF<int>(buf);
+    o_w = readBUF<int>(buf);
+    o_n = readBUF<int>(buf);
+    p = readBUF<float>(buf);
+    assert(buf == bufCheck + length);
+}
+
+GeneralizedMeanPoolingPRT::~GeneralizedMeanPoolingPRT() {
+
+}
+
+int GeneralizedMeanPoolingPRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims GeneralizedMeanPoolingPRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return Dims3{this->o_c, this->o_h, this->o_w};
+}
+
+int GeneralizedMeanPoolingPRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+void GeneralizedMeanPoolingPRT::terminate() NOEXCEPT {
+
+}
+
+size_t GeneralizedMeanPoolingPRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
+    return 0;
+}
+
+#if NV_TENSORRT_MAJOR > 7
+int GeneralizedMeanPoolingPRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                                cudaStream_t stream) NOEXCEPT {
+    dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
+    dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
+    GeneralizedMeanPoolingP(srcData, dstData, batchSize, this->i_c, this->i_h, this->i_w, this->p, stream);
+    return 0;
+}
+#elif NV_TENSORRT_MAJOR <= 7
+int32_t GeneralizedMeanPoolingPRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace,
+                                    cudaStream_t stream) {
+    dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
+    dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
+    GeneralizedMeanPoolingP(srcData, dstData, batchSize, this->i_c, this->i_h, this->i_w, this->p, stream);
+    return 0;
+}
+#endif
+
+
+size_t GeneralizedMeanPoolingPRT::getSerializationSize() const NOEXCEPT {
+    return 8*sizeof(int) + sizeof(float);
+}
+
+void GeneralizedMeanPoolingPRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char*>(buffer),*a=buf;
+    writeBUF(buf, this->i_c);
+    writeBUF(buf, this->i_h);
+    writeBUF(buf, this->i_w);
+    writeBUF(buf, this->i_n);
+    writeBUF(buf, this->o_c);
+    writeBUF(buf, this->o_h);
+    writeBUF(buf, this->o_w);
+    writeBUF(buf, this->o_n);
+    writeBUF(buf, this->p);
+    assert(buf == a + getSerializationSize());
+}
+
+void GeneralizedMeanPoolingPRT::destroy() NOEXCEPT {
+delete this;
+}
+
+bool GeneralizedMeanPoolingPRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR);
+}
+
+const char *GeneralizedMeanPoolingPRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void GeneralizedMeanPoolingPRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *GeneralizedMeanPoolingPRT::getPluginType() const NOEXCEPT {
+    return GENERALIZEDMEANPOOLINGPRT_PLUGIN_NAME;
+}
+
+const char *GeneralizedMeanPoolingPRT::getPluginVersion() const NOEXCEPT {
+    return GENERALIZEDMEANPOOLINGPRT_PLUGIN_VERSION;
+}
+
+IPluginV2Ext *GeneralizedMeanPoolingPRT::clone() const NOEXCEPT {
+    auto *pl = new GeneralizedMeanPoolingPRT(i_c,i_h,i_w,i_n, o_c,o_h,o_w,o_n,p);
+    pl->setPluginNamespace(mPluginNamespace.c_str());
+    return pl;
+}
+
+DataType
+GeneralizedMeanPoolingPRT::getOutputDataType(int index, const nvinfer1::DataType *inputTypes, int nbInputs) const NOEXCEPT {
+    return DataType::kFLOAT;
+}
+
+void GeneralizedMeanPoolingPRT::attachToContext(cudnnContext *cudnnContext, cublasContext *cublasContext,
+                                         IGpuAllocator *gpuAllocator) NOEXCEPT {
+
+}
+
+bool GeneralizedMeanPoolingPRT::isOutputBroadcastAcrossBatch(int outputIndex, const bool *inputIsBroadcasted,
+                                                      int nbInputs) const NOEXCEPT {
+    return false;
+}
+
+bool GeneralizedMeanPoolingPRT::canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT {
+    return false;
+}
+
+void
+GeneralizedMeanPoolingPRT::configurePlugin(const Dims *inputDims, int32_t nbInputs, const Dims *outputDims, int32_t nbOutputs,
+                                    const DataType *inputTypes, const DataType *outputTypes,
+                                    const bool *inputIsBroadcast, const bool *outputIsBroadcast,
+                                    PluginFormat floatFormat, int32_t maxBatchSize) NOEXCEPT {
+
+}
+
+void GeneralizedMeanPoolingPRT::detachFromContext() NOEXCEPT {
+    IPluginV2Ext::detachFromContext();
+}
+
+GeneralizedMeanPoolingPRTPluginCreator::GeneralizedMeanPoolingPRTPluginCreator() {
+    mPluginAttributes.clear();
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void GeneralizedMeanPoolingPRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *GeneralizedMeanPoolingPRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2Ext *GeneralizedMeanPoolingPRTPluginCreator::deserializePlugin(const char *name, const void *serialData,size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new GeneralizedMeanPoolingPRT(serialData,serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2Ext *GeneralizedMeanPoolingPRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    const PluginField *fields = fc->fields;
+    int i_c = *(static_cast<const int *>(fields[0].data));
+    int i_h = *(static_cast<const int *>(fields[1].data));
+    int i_w = *(static_cast<const int *>(fields[2].data));
+    int i_n = *(static_cast<const int *>(fields[3].data));
+    int o_c = *(static_cast<const int *>(fields[4].data));
+    int o_h = *(static_cast<const int *>(fields[5].data));
+    int o_w = *(static_cast<const int *>(fields[6].data));
+    int o_n = *(static_cast<const int *>(fields[7].data));
+    float p = *(static_cast<const float *>(fields[8].data));
+    auto *pluginObj = new GeneralizedMeanPoolingPRT(i_c,i_h,i_w,i_n, o_c,o_h,o_w,o_n,p);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *GeneralizedMeanPoolingPRTPluginCreator::getPluginName() const NOEXCEPT {
+    return GENERALIZEDMEANPOOLINGPRT_PLUGIN_NAME;
+}
+
+const char *GeneralizedMeanPoolingPRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return GENERALIZEDMEANPOOLINGPRT_PLUGIN_VERSION;
+}
+
+const PluginFieldCollection *GeneralizedMeanPoolingPRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+
+
+
+
+
+
+
+
+
+
+

tests/backbones/resnet101/Resnet-101-AP-GeM.cpp

@@ -0,0 +1,343 @@
+#include <iostream>
+#include "tkDNN/tkdnn.h"
+
+const char *input_bin = "Resnet-101-AP-GeM/debug/input.bin";
+const char *conv1_bin = "Resnet-101-AP-GeM/layers/module-conv1.bin";
+
+//layer1
+const char *layer1_bin[]={
+"Resnet-101-AP-GeM/layers/module-layer1-0-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer1-0-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer1-0-conv3.bin",
+"Resnet-101-AP-GeM/layers/module-layer1-0-downsample-0.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer1-1-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer1-1-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer1-1-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer1-2-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer1-2-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer1-2-conv3.bin"};
+                        
+
+//layer2
+const char *layer2_bin[]={
+"Resnet-101-AP-GeM/layers/module-layer2-0-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer2-0-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer2-0-conv3.bin",
+"Resnet-101-AP-GeM/layers/module-layer2-0-downsample-0.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer2-1-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer2-1-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer2-1-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer2-2-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer2-2-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer2-2-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer2-3-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer2-3-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer2-3-conv3.bin"
+};
+//layer3
+const char *layer3_bin[]={
+"Resnet-101-AP-GeM/layers/module-layer3-0-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-0-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-0-conv3.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-0-downsample-0.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-1-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-1-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-1-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-2-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-2-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-2-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-3-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-3-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-3-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-4-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-4-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-4-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-5-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-5-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-5-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-6-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-6-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-6-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-7-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-7-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-7-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-8-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-8-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-8-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-9-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-9-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-9-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-10-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-10-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-10-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-11-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-11-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-11-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-12-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-12-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-12-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-13-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-13-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-13-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-14-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-14-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-14-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-15-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-15-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-15-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-16-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-16-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-16-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-17-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-17-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-17-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-18-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-18-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-18-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-19-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-19-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-19-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-20-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-20-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-20-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-21-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-21-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-21-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer3-22-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-22-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer3-22-conv3.bin"};
+
+
+//layer4
+const char *layer4_bin[]={
+"Resnet-101-AP-GeM/layers/module-layer4-0-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer4-0-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer4-0-conv3.bin",
+"Resnet-101-AP-GeM/layers/module-layer4-0-downsample-0.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer4-1-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer4-1-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer4-1-conv3.bin",
+
+"Resnet-101-AP-GeM/layers/module-layer4-2-conv1.bin",
+"Resnet-101-AP-GeM/layers/module-layer4-2-conv2.bin",
+"Resnet-101-AP-GeM/layers/module-layer4-2-conv3.bin"};
+
+//final
+const char *fc_bin = "Resnet-101-AP-GeM/layers/module-fc.bin";
+
+// const char *output_bin = "Resnet-101-AP-GeM/debug/module-layer4-2-bn3.bin";
+// const char *output_bin = "Resnet-101-AP-GeM/debug/module-adpool.bin";
+const char *output_bin = "Resnet-101-AP-GeM/debug/module-fc.bin";
+
+int main()
+{
+
+    std::string bin_path  = "Resnet-101-AP-GeM";
+    // downloadWeightsifDoNotExist(input_bin, bin_path, "https://cloud.hipert.unimore.it/s/2MyPcWnEzGTm28A/download");
+
+    // Network layout
+    tk::dnn::dataDim_t dim(1, 3, 768, 1024, 1);
+    tk::dnn::Network net(dim);
+
+    tk::dnn::Conv2d conv1(&net, 64, 7, 7, 2, 2, 3, 3, conv1_bin, true);
+    tk::dnn::Activation relu3(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Pooling maxpool4(&net, 3, 3, 2, 2, 1, 1, tk::dnn::POOLING_MAX);
+
+    //layer 1
+    int id_layer1_bin = 0;
+    tk::dnn::Layer *last = &maxpool4;
+    for(int i=0; i<3;i++)
+    {
+        tk::dnn::Conv2d     *layer1_0_conv1 = new tk::dnn::Conv2d(&net, 64, 1, 1, 1, 1, 0, 0, layer1_bin[id_layer1_bin++], true);
+        tk::dnn::Activation *relu1_0_1      = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU);
+        tk::dnn::Conv2d     *layer1_0_conv2 = new tk::dnn::Conv2d(&net, 64, 3, 3, 1, 1, 1, 1, layer1_bin[id_layer1_bin++], true);
+        tk::dnn::Activation *relu1_0_2      = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU);
+        tk::dnn::Conv2d     *layer1_0_conv3 = new tk::dnn::Conv2d(&net, 256, 1, 1, 1, 1, 0, 0, layer1_bin[id_layer1_bin++], true);
+        if(i==0) {
+            tk::dnn::Layer    *route_1_0_layers[1] = { last };
+            tk::dnn::Route    *route_1_0             = new tk::dnn::Route(&net, route_1_0_layers, 1);
+            tk::dnn::Conv2d   *layer1_0_downsample_0 = new tk::dnn::Conv2d(&net, 256, 1, 1, 1, 1, 0, 0, layer1_bin[id_layer1_bin++], true);
+            tk::dnn::Shortcut *s1_0                  = new tk::dnn::Shortcut(&net, layer1_0_conv3);
+        } else {
+            tk::dnn::Shortcut *s1_0                  = new tk::dnn::Shortcut(&net, last);
+        }
+        tk::dnn::Activation   *layer1_0_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU);
+        last = layer1_0_relu;
+    }
+
+    // tk::dnn::Activation *last_activation = (tk::dnn::Activation *) net.layers[net.num_layers-1];
+    // layer 2
+    int id_layer2_bin = 0;
+    for(int i=0; i<4;i++)
+    {
+        tk::dnn::Conv2d     *layer1_0_conv1 = new tk::dnn::Conv2d(&net, 128, 1, 1, 1, 1, 0, 0, layer2_bin[id_layer2_bin++], true);
+        tk::dnn::Activation *relu1_0_1      = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU);
+        tk::dnn::Conv2d     *layer1_0_conv2;
+        if(i==0)
+            layer1_0_conv2 = new tk::dnn::Conv2d(&net, 128, 3, 3, 2, 2, 1, 1, layer2_bin[id_layer2_bin++], true);
+        else
+            layer1_0_conv2 = new tk::dnn::Conv2d(&net, 128, 3, 3, 1, 1, 1, 1, layer2_bin[id_layer2_bin++], true);
+        
+        tk::dnn::Activation *relu1_0_2      = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU);
+        tk::dnn::Conv2d     *layer1_0_conv3 = new tk::dnn::Conv2d(&net, 512, 1, 1, 1, 1, 0, 0, layer2_bin[id_layer2_bin++], true);
+        if(i==0) 
+        {
+            tk::dnn::Layer    *route_1_0_layers[1] = { last };
+            tk::dnn::Route    *route_1_0             = new tk::dnn::Route(&net, route_1_0_layers, 1);
+            tk::dnn::Conv2d   *layer1_0_downsample_0 = new tk::dnn::Conv2d(&net, 512, 1, 1, 2, 2, 0, 0, layer2_bin[id_layer2_bin++], true);
+            tk::dnn::Shortcut *s1_0                  = new tk::dnn::Shortcut(&net, layer1_0_conv3);
+        } 
+        else 
+        {
+            tk::dnn::Shortcut *s1_0                  = new tk::dnn::Shortcut(&net, last);
+        }
+        tk::dnn::Activation   *layer1_0_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU);
+        last = layer1_0_relu;
+    }
+    
+    // layer 3
+    int id_layer3_bin = 0;
+    for(int i=0; i<23;i++)
+    {
+        tk::dnn::Conv2d     *layer1_0_conv1 = new tk::dnn::Conv2d(&net, 256, 1, 1, 1, 1, 0, 0, layer3_bin[id_layer3_bin++], true);
+        tk::dnn::Activation *relu1_0_1      = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU);
+        tk::dnn::Conv2d     *layer1_0_conv2;
+        if(i==0)
+            layer1_0_conv2 = new tk::dnn::Conv2d(&net, 256, 3, 3, 2, 2, 1, 1, layer3_bin[id_layer3_bin++], true);
+        else
+            layer1_0_conv2 = new tk::dnn::Conv2d(&net, 256, 3, 3, 1, 1, 1, 1, layer3_bin[id_layer3_bin++], true);
+        
+        tk::dnn::Activation *relu1_0_2      = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU);
+        tk::dnn::Conv2d     *layer1_0_conv3 = new tk::dnn::Conv2d(&net, 1024, 1, 1, 1, 1, 0, 0, layer3_bin[id_layer3_bin++], true);
+        if(i==0) 
+        {
+            tk::dnn::Layer    *route_1_0_layers[1] = { last };
+            tk::dnn::Route    *route_1_0             = new tk::dnn::Route(&net, route_1_0_layers, 1);
+            tk::dnn::Conv2d   *layer1_0_downsample_0 = new tk::dnn::Conv2d(&net, 1024, 1, 1, 2, 2, 0, 0, layer3_bin[id_layer3_bin++], true);
+            tk::dnn::Shortcut *s1_0                  = new tk::dnn::Shortcut(&net, layer1_0_conv3);
+        } 
+        else 
+        {
+            tk::dnn::Shortcut *s1_0                  = new tk::dnn::Shortcut(&net, last);
+        }
+        tk::dnn::Activation   *layer1_0_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU);
+        last = layer1_0_relu;
+    }
+
+    // layer 4
+    int id_layer4_bin = 0;
+    for(int i=0; i<3;i++)
+    {
+        tk::dnn::Conv2d     *layer1_0_conv1 = new tk::dnn::Conv2d(&net, 512, 1, 1, 1, 1, 0, 0, layer4_bin[id_layer4_bin++], true);
+        tk::dnn::Activation *relu1_0_1      = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU);
+        tk::dnn::Conv2d     *layer1_0_conv2;
+        if(i==0)
+            layer1_0_conv2 = new tk::dnn::Conv2d(&net, 512, 3, 3, 2, 2, 1, 1, layer4_bin[id_layer4_bin++], true);
+        else
+            layer1_0_conv2 = new tk::dnn::Conv2d(&net, 512, 3, 3, 1, 1, 1, 1, layer4_bin[id_layer4_bin++], true);
+        
+        tk::dnn::Activation *relu1_0_2      = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU);
+        tk::dnn::Conv2d     *layer1_0_conv3 = new tk::dnn::Conv2d(&net, 2048, 1, 1, 1, 1, 0, 0, layer4_bin[id_layer4_bin++], true);
+        if(i==0) 
+        {
+            tk::dnn::Layer    *route_1_0_layers[1] = { last };
+            tk::dnn::Route    *route_1_0             = new tk::dnn::Route(&net, route_1_0_layers, 1);
+            tk::dnn::Conv2d   *layer1_0_downsample_0 = new tk::dnn::Conv2d(&net, 2048, 1, 1, 2, 2, 0, 0, layer4_bin[id_layer4_bin++], true);
+            tk::dnn::Shortcut *s1_0                  = new tk::dnn::Shortcut(&net, layer1_0_conv3);
+        } 
+        else 
+        {
+            tk::dnn::Shortcut *s1_0                  = new tk::dnn::Shortcut(&net, last);
+        }
+        tk::dnn::Activation   *layer1_0_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU);
+        if(i == 2)
+            layer1_0_relu->setFinal();
+        last = layer1_0_relu;
+    }
+
+    //final
+    tk::dnn::Pooling avgpool(&net, 24, 32, 24, 32, 0, 0, tk::dnn::POOLING_GENERALIZED_MEAN_P, 2.9963f);
+    tk::dnn::Dense   fc(&net, 2048, fc_bin);
+
+    // Load input
+    dnnType *data;
+    dnnType *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+    //printDeviceVector(64, data, true);
+
+    //print network model
+    net.print();
+
+    //convert network to tensorRT
+    tk::dnn::NetworkRT netRT(&net, net.getNetworkRTName("Resnet-101-AP-GeM"));
+
+    
+    tk::dnn::dataDim_t out_dim;
+    out_dim = net.layers[net.num_layers-1]->output_dim;
+    dnnType *cudnn_out, *rt_out;
+
+    tk::dnn::dataDim_t dim1 = dim; //input dim
+    printCenteredTitle(" CUDNN inference ", '=', 30);
+    {
+        dim1.print();
+        TKDNN_TSTART
+        net.infer(dim1, data);
+        TKDNN_TSTOP
+        dim1.print();
+    }
+    cudnn_out = net.layers[net.num_layers-1]->dstData;
+
+    tk::dnn::dataDim_t dim2 = dim;
+    printCenteredTitle(" TENSORRT inference ", '=', 30);
+    {
+        dim2.print();
+        TKDNN_TSTART
+        netRT.infer(dim2, data);
+        TKDNN_TSTOP
+        dim2.print();
+    }
+    rt_out = (dnnType *)netRT.buffersRT[2];
+
+
+    printCenteredTitle(std::string(" RESNET CHECK RESULTS ").c_str(), '=', 30);
+    dnnType *out, *out_h;
+    int odim = out_dim.tot();
+    readBinaryFile(output_bin, odim, &out_h, &out);
+    
+    std::cout<<"CUDNN vs correct"; 
+    int ret_cudnn = checkResult(odim, cudnn_out, out) == 0 ? 0: ERROR_CUDNN;
+    std::cout<<"TRT   vs correct"; 
+    int ret_tensorrt = checkResult(odim, rt_out, out) == 0 ? 0 : ERROR_TENSORRT;
+    std::cout<<"CUDNN vs TRT    "; 
+    int ret_cudnn_tensorrt = checkResult(odim, cudnn_out, rt_out) == 0 ? 0 : ERROR_CUDNNvsTENSORRT;
+
+    return ret_cudnn | ret_tensorrt | ret_cudnn_tensorrt;
+}