diff --git a/CMakeLists.txt b/CMakeLists.txt index c29e987..2b6dd5c 100644 --- a/CMakeLists.txt +++ b/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) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index daa27e5..ba0377e 100644 --- a/include/tkDNN/Layer.h +++ b/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; }; diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index 4d5474b..ad8beb9 100644 --- a/include/tkDNN/kernels.h +++ b/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)); diff --git a/include/tkDNN/pluginsRT/GeneralizedMeanPoolingPRT.h b/include/tkDNN/pluginsRT/GeneralizedMeanPoolingPRT.h new file mode 100644 index 0000000..0226335 --- /dev/null +++ b/include/tkDNN/pluginsRT/GeneralizedMeanPoolingPRT.h @@ -0,0 +1,103 @@ +#include +#include "../kernels.h" +#include +#include +#include + + +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 mPluginAttributes; + std::string mPluginNamespace; + + }; + + REGISTER_TENSORRT_PLUGIN(GeneralizedMeanPoolingPRTPluginCreator); +}; diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 26489bf..8484fb5 100644 --- a/src/NetworkRT.cpp +++ b/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 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 { diff --git a/src/Pooling.cpp b/src/Pooling.cpp index 2cea2f6..d44c815 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) : + 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); diff --git a/src/kernels/pooling.cu b/src/kernels/pooling.cu index cd80a6d..d68bf35 100644 --- a/src/kernels/pooling.cu +++ b/src/kernels/pooling.cu @@ -39,6 +39,47 @@ __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; + + std::cerr<<"Calling forward_gen_avgpool_p_layer_kernel "<>>(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) { diff --git a/src/pluginsRT/GeneralizedMeanPoolingPRT.cpp b/src/pluginsRT/GeneralizedMeanPoolingPRT.cpp new file mode 100644 index 0000000..5509f1d --- /dev/null +++ b/src/pluginsRT/GeneralizedMeanPoolingPRT.cpp @@ -0,0 +1,216 @@ +#include +using namespace nvinfer1; + +std::vector 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(data),*bufCheck = buf; + i_c = readBUF(buf); + i_h = readBUF(buf); + i_w = readBUF(buf); + i_n = readBUF(buf); + o_c = readBUF(buf); + o_h = readBUF(buf); + o_w = readBUF(buf); + o_n = readBUF(buf); + p = readBUF(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(inputs[0]); + dnnType *dstData = reinterpret_cast(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(inputs[0]); + dnnType *dstData = reinterpret_cast(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); +} + +void GeneralizedMeanPoolingPRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(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(fields[0].data)); + int i_h = *(static_cast(fields[1].data)); + int i_w = *(static_cast(fields[2].data)); + int i_n = *(static_cast(fields[3].data)); + int o_c = *(static_cast(fields[4].data)); + int o_h = *(static_cast(fields[5].data)); + int o_w = *(static_cast(fields[6].data)); + int o_n = *(static_cast(fields[7].data)); + float p = *(static_cast(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; +} + + + + + + + + + + + + diff --git a/tests/backbones/resnet101/Resnet-101-AP-GeM.cpp b/tests/backbones/resnet101/Resnet-101-AP-GeM.cpp new file mode 100644 index 0000000..ed86052 --- /dev/null +++ b/tests/backbones/resnet101/Resnet-101-AP-GeM.cpp @@ -0,0 +1,343 @@ +#include +#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[1]; + + + 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; +}