From f58776c046d5792fb37966c3a7f80384239a7bfa Mon Sep 17 00:00:00 2001 From: perseusdg Date: Mon, 16 Aug 2021 14:04:17 +0530 Subject: [PATCH 1/6] Revert "fix return type for bool functions" This reverts commit d2ce966c46877495a1d9fa83edab45620173ec4e. Undo tensorrt8 changes on the master branch --- include/tkDNN/SegmentationNN.h | 1 - src/CenterTrack.cpp | 4 +--- src/CenternetDetection.cpp | 1 - src/CenternetDetection3D.cpp | 1 - 4 files changed, 1 insertion(+), 6 deletions(-) diff --git a/include/tkDNN/SegmentationNN.h b/include/tkDNN/SegmentationNN.h index 403bc28..ad21d74 100644 --- a/include/tkDNN/SegmentationNN.h +++ b/include/tkDNN/SegmentationNN.h @@ -183,7 +183,6 @@ class SegmentationNN { checkCuda(cudaMemcpyAsync(mean_d, mean.data(), mean.size() * sizeof(float), cudaMemcpyHostToDevice, netRT->stream)); checkCuda(cudaMemcpyAsync(stddev_d, stddev.data(), stddev.size() * sizeof(float), cudaMemcpyHostToDevice, netRT->stream)); - return true; } diff --git a/src/CenterTrack.cpp b/src/CenterTrack.cpp index 55bf82f..7931488 100644 --- a/src/CenterTrack.cpp +++ b/src/CenterTrack.cpp @@ -60,7 +60,6 @@ bool CenterTrack::init_preprocessing(){ checkCuda( cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot() * nBatches)); checkCuda( cudaMalloc(&input_pre_inf_d, sizeof(dnnType)*dim.tot())); checkCuda( cudaMalloc(&d_ptrs, dim.tot() * sizeof(float)) ); - return true; } bool CenterTrack::init_pre_inf(){ @@ -204,7 +203,6 @@ bool CenterTrack::init_postprocessing(){ trRes.resize(nBatches); countTr.resize(nBatches, 0); trackId.resize(nBatches, 0); - return true; } bool CenterTrack::init_visualization(const int n_classes){ @@ -313,7 +311,7 @@ void CenterTrack::preprocess(cv::Mat &frame, const int bi){ } float c[] = {new_width / 2.0f, new_height /2.0f}; - float s[] = {static_cast(dim.w), static_cast(dim.h)}; + float s[] = {dim.w, dim.h}; // float s = new_width >= new_height ? new_width : new_height; // ----------- get_affine_transform // rot_rad = pi * 0 / 100 --> 0 diff --git a/src/CenternetDetection.cpp b/src/CenternetDetection.cpp index a06bc11..bd09005 100644 --- a/src/CenternetDetection.cpp +++ b/src/CenternetDetection.cpp @@ -118,7 +118,6 @@ bool CenternetDetection::init(const std::string& tensor_path, const int n_classe dst2.at(2,0)=dst2.at(1,0) + (-dst2.at(0,1)+dst2.at(1,1) ); dst2.at(2,1)=dst2.at(1,1) + (dst2.at(0,0)-dst2.at(1,0) ); - return true; } diff --git a/src/CenternetDetection3D.cpp b/src/CenternetDetection3D.cpp index fd09132..653624b 100644 --- a/src/CenternetDetection3D.cpp +++ b/src/CenternetDetection3D.cpp @@ -167,7 +167,6 @@ bool CenternetDetection3D::init(const std::string& tensor_path, const int n_clas faceId.push_back({2,3,7,6}); faceId.push_back({3,0,4,7}); // ([[0,1,5,4], [1,2,6, 5], [2,3,7,6], [3,0,4,7]]); - return true; } void CenternetDetection3D::preprocess(cv::Mat &frame, const int bi){ From 52db97a140767a2063437897bfc40761216a7e9d Mon Sep 17 00:00:00 2001 From: perseusdg Date: Mon, 16 Aug 2021 14:06:56 +0530 Subject: [PATCH 2/6] Revert "Revert "fix return type for bool functions"" This reverts commit f58776c046d5792fb37966c3a7f80384239a7bfa. --- include/tkDNN/SegmentationNN.h | 1 + src/CenterTrack.cpp | 4 +++- src/CenternetDetection.cpp | 1 + src/CenternetDetection3D.cpp | 1 + 4 files changed, 6 insertions(+), 1 deletion(-) diff --git a/include/tkDNN/SegmentationNN.h b/include/tkDNN/SegmentationNN.h index ad21d74..403bc28 100644 --- a/include/tkDNN/SegmentationNN.h +++ b/include/tkDNN/SegmentationNN.h @@ -183,6 +183,7 @@ class SegmentationNN { checkCuda(cudaMemcpyAsync(mean_d, mean.data(), mean.size() * sizeof(float), cudaMemcpyHostToDevice, netRT->stream)); checkCuda(cudaMemcpyAsync(stddev_d, stddev.data(), stddev.size() * sizeof(float), cudaMemcpyHostToDevice, netRT->stream)); + return true; } diff --git a/src/CenterTrack.cpp b/src/CenterTrack.cpp index 7931488..55bf82f 100644 --- a/src/CenterTrack.cpp +++ b/src/CenterTrack.cpp @@ -60,6 +60,7 @@ bool CenterTrack::init_preprocessing(){ checkCuda( cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot() * nBatches)); checkCuda( cudaMalloc(&input_pre_inf_d, sizeof(dnnType)*dim.tot())); checkCuda( cudaMalloc(&d_ptrs, dim.tot() * sizeof(float)) ); + return true; } bool CenterTrack::init_pre_inf(){ @@ -203,6 +204,7 @@ bool CenterTrack::init_postprocessing(){ trRes.resize(nBatches); countTr.resize(nBatches, 0); trackId.resize(nBatches, 0); + return true; } bool CenterTrack::init_visualization(const int n_classes){ @@ -311,7 +313,7 @@ void CenterTrack::preprocess(cv::Mat &frame, const int bi){ } float c[] = {new_width / 2.0f, new_height /2.0f}; - float s[] = {dim.w, dim.h}; + float s[] = {static_cast(dim.w), static_cast(dim.h)}; // float s = new_width >= new_height ? new_width : new_height; // ----------- get_affine_transform // rot_rad = pi * 0 / 100 --> 0 diff --git a/src/CenternetDetection.cpp b/src/CenternetDetection.cpp index bd09005..a06bc11 100644 --- a/src/CenternetDetection.cpp +++ b/src/CenternetDetection.cpp @@ -118,6 +118,7 @@ bool CenternetDetection::init(const std::string& tensor_path, const int n_classe dst2.at(2,0)=dst2.at(1,0) + (-dst2.at(0,1)+dst2.at(1,1) ); dst2.at(2,1)=dst2.at(1,1) + (dst2.at(0,0)-dst2.at(1,0) ); + return true; } diff --git a/src/CenternetDetection3D.cpp b/src/CenternetDetection3D.cpp index 653624b..fd09132 100644 --- a/src/CenternetDetection3D.cpp +++ b/src/CenternetDetection3D.cpp @@ -167,6 +167,7 @@ bool CenternetDetection3D::init(const std::string& tensor_path, const int n_clas faceId.push_back({2,3,7,6}); faceId.push_back({3,0,4,7}); // ([[0,1,5,4], [1,2,6, 5], [2,3,7,6], [3,0,4,7]]); + return true; } void CenternetDetection3D::preprocess(cv::Mat &frame, const int bi){ From 19b1d6d6f7c1fa1db7bacc45ef50537fb24949b3 Mon Sep 17 00:00:00 2001 From: perseusdg Date: Mon, 16 Aug 2021 14:07:07 +0530 Subject: [PATCH 3/6] Revert "Updates to build libkernel.so under TensorRT 8" This reverts commit eba78e7e78093e534938b2dc32a3cfd2a16ee9e4. --- CMakeLists.txt | 2 +- include/tkDNN/Int8BatchStream.h | 4 +- include/tkDNN/Int8Calibrator.h | 8 +- include/tkDNN/NetworkRT.h | 7 +- include/tkDNN/pluginsRT/ActivationLeakyRT.h | 168 +++------ .../tkDNN/pluginsRT/ActivationLogisticRT.h | 107 +----- include/tkDNN/pluginsRT/ActivationMishRT.h | 159 +++----- .../tkDNN/pluginsRT/ActivationReLUCeilingRT.h | 172 +++------ include/tkDNN/pluginsRT/DeformableConvRT.h | 203 +--------- include/tkDNN/pluginsRT/FlattenConcatRT.h | 119 +----- .../tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h | 139 +------ include/tkDNN/pluginsRT/RegionRT.h | 125 +------ include/tkDNN/pluginsRT/ReorgRT.h | 115 +----- include/tkDNN/pluginsRT/ReshapeRT.h | 122 +----- include/tkDNN/pluginsRT/ResizeLayerRT.h | 129 +------ include/tkDNN/pluginsRT/RouteRT.h | 127 ++----- include/tkDNN/pluginsRT/ShortcutRT.h | 137 ++----- include/tkDNN/pluginsRT/UpsampleRT.h | 123 ++----- include/tkDNN/pluginsRT/YoloRT.h | 347 ++++++------------ include/tkDNN/utils.h | 10 - src/Int8Calibrator.cpp | 2 +- 21 files changed, 447 insertions(+), 1878 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 3e823dc..b366416 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -3,7 +3,7 @@ cmake_minimum_required(VERSION 3.15) project (tkDNN) set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake) if(UNIX) -set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++14 -fPIC -Wno-deprecated-declarations -Wno-unused-variable ") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC -Wno-deprecated-declarations -Wno-unused-variable ") endif() if(WIN32) set(CMAKE_CXX_STANDARD 11) diff --git a/include/tkDNN/Int8BatchStream.h b/include/tkDNN/Int8BatchStream.h index 028f75d..c39a11c 100644 --- a/include/tkDNN/Int8BatchStream.h +++ b/include/tkDNN/Int8BatchStream.h @@ -39,7 +39,7 @@ public: float *getLabels() { return mLabels.data(); } int getBatchesRead() const { return mBatchCount; } int getBatchSize() const { return mBatchSize; } - nvinfer1::Dims4 getDims() const { return mDims; } + nvinfer1::DimsNCHW getDims() const { return mDims; } float* getFileBatch() { return &mFileBatch[0]; } float* getFileLabels() { return &mFileLabels[0]; } void readInListFile(const std::string& dataFilePath, std::vector& mListIn); @@ -55,7 +55,7 @@ private: int mFileBatchPos{ 0 }; int mImageSize{ 0 }; - nvinfer1::Dims4 mDims; + nvinfer1::DimsNCHW mDims; std::vector mBatch; std::vector mLabels; std::vector mFileBatch; diff --git a/include/tkDNN/Int8Calibrator.h b/include/tkDNN/Int8Calibrator.h index c7a7d97..4a0ea47 100644 --- a/include/tkDNN/Int8Calibrator.h +++ b/include/tkDNN/Int8Calibrator.h @@ -30,10 +30,10 @@ public: Int8EntropyCalibrator(BatchStream& stream, int firstBatch, const std::string& calibTableFilePath, const std::string& inputBlobName, bool readCache = true); virtual ~Int8EntropyCalibrator() { checkCuda(cudaFree(mDeviceInput)); } - int getBatchSize() const NOEXCEPT override { return mStream.getBatchSize(); } - bool getBatch(void* bindings[], const char* names[], int nbBindings) NOEXCEPT override; - const void* readCalibrationCache(size_t& length) NOEXCEPT override; - void writeCalibrationCache(const void* cache, size_t length) NOEXCEPT override; + int getBatchSize() const override { return mStream.getBatchSize(); } + bool getBatch(void* bindings[], const char* names[], int nbBindings) override; + const void* readCalibrationCache(size_t& length) override; + void writeCalibrationCache(const void* cache, size_t length) override; private: BatchStream mStream; diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 53d8d91..9892a24 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -40,16 +40,14 @@ using namespace nvinfer1; #include "pluginsRT/ReshapeRT.h" #include "pluginsRT/MaxPoolingFixedSizeRT.h" -/* -class PluginFactory : IPlugin +class PluginFactory : IPluginFactory { public: YoloRT *yolos[16]; int n_yolos; virtual IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength); -};*/ - +}; @@ -76,6 +74,7 @@ public: dnnType *output; cudaStream_t stream; + PluginFactory *pluginFactory; NetworkRT(Network *net, const char *name); virtual ~NetworkRT(); diff --git a/include/tkDNN/pluginsRT/ActivationLeakyRT.h b/include/tkDNN/pluginsRT/ActivationLeakyRT.h index 6005013..330ed37 100644 --- a/include/tkDNN/pluginsRT/ActivationLeakyRT.h +++ b/include/tkDNN/pluginsRT/ActivationLeakyRT.h @@ -1,147 +1,61 @@ +#include #include "../kernels.h" -#include -class ActivationLeakyRT : public IPluginV2 { +class ActivationLeakyRT : public IPlugin { public: - ActivationLeakyRT(float s) { slope = s; } + ActivationLeakyRT(float s) { + slope = s; + } - ActivationLeakyRT(const void *data, size_t length) - { - const char* buf = reinterpret_cast(data),*bufCheck = buf; - slope = readBUF(buf); - size = readBUF(buf); - assert(buf == bufCheck + length); + ~ActivationLeakyRT(){ - } + } - ~ActivationLeakyRT() {} + int getNbOutputs() const override { + return 1; + } - int getNbOutputs() const NOEXCEPT override { return 1; } + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return inputs[0]; + } - Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override { - return inputs[0]; - } + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + size = 1; + for(int i=0; i(inputs[0]), - reinterpret_cast(outputs[0]), batchSize * size, slope, - stream); - return 0; - } + virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { - virtual size_t getSerializationSize() const NOEXCEPT override { - return 1 * sizeof(int) + 1 * sizeof(float); - } + activationLEAKYForward((dnnType*)reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize*size, slope, stream); + return 0; + } - virtual void serialize(void *buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer), *a = buf; - tk::dnn::writeBUF(buf, size); - assert(buf == a + getSerializationSize()); - } - bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override { - return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR); - } + virtual size_t getSerializationSize() override { + return 1*sizeof(int) + 1*sizeof(float); + } - const char *getPluginType() const NOEXCEPT override { - return "ActivationLeakyRT_tkDNN"; - } + virtual void serialize(void* buffer) override { + char *buf = reinterpret_cast(buffer),*a=buf; + tk::dnn::writeBUF(buf, size); + assert(buf == a + getSerializationSize()); + } - const char *getPluginVersion() const NOEXCEPT override { - return "1"; - } - - void destroy() NOEXCEPT override { delete this; } - - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override { - mPluginNamespace = pluginNamespace; - } - - IPluginV2* clone() const NOEXCEPT override { - ActivationLeakyRT *p = new ActivationLeakyRT(slope); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - - int size; - float slope; - -private: - std::string mPluginNamespace; + int size; + float slope; }; - -class ActivationLeakyRTPluginCreator : public IPluginCreator { -public: - ActivationLeakyRTPluginCreator() { - mPluginAttributes.emplace_back( - PluginField("slope", nullptr, PluginFieldType::kFLOAT32, 1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override { - ActivationLeakyRT *pluginObj = new ActivationLeakyRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } - - IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override { - const PluginField *fields = fc->fields; - assert(fc->nbFields == 1); - assert(fields[0].type == PluginFieldType::kFLOAT32); - float slope = *(static_cast(fields[0].data)); - ActivationLeakyRT *pluginObj = new ActivationLeakyRT(slope); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ActivationLeakyRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ActivationLeakyRTPluginCreator); \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ActivationLogisticRT.h b/include/tkDNN/pluginsRT/ActivationLogisticRT.h index bf4bcfd..063931f 100644 --- a/include/tkDNN/pluginsRT/ActivationLogisticRT.h +++ b/include/tkDNN/pluginsRT/ActivationLogisticRT.h @@ -1,18 +1,11 @@ #include #include "../kernels.h" -class ActivationLogisticRT : public IPluginV2 { +class ActivationLogisticRT : public IPlugin { public: ActivationLogisticRT() { - } - - ActivationLogisticRT(const void *data, size_t length) - { - const char* buf = reinterpret_cast(data),*bufCheck = buf; - size = readBUF(buf); - assert(buf == bufCheck + length); } @@ -20,33 +13,33 @@ public: } - int getNbOutputs() const NOEXCEPT override { + int getNbOutputs() const override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { return inputs[0]; } - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format, int maxBatchSize) NOEXCEPT override { + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { size = 1; for(int i=0; i(inputs[0]), reinterpret_cast(outputs[0]), batchSize*size, stream); @@ -54,94 +47,14 @@ public: } - virtual size_t getSerializationSize() const NOEXCEPT override { + virtual size_t getSerializationSize() override { return 1*sizeof(int); } - virtual void serialize(void* buffer) const NOEXCEPT override { + virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer); tk::dnn::writeBUF(buf, size); } - const char *getPluginType() const NOEXCEPT override { - return "ActivationLogisticRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override { - return "1"; - } - - void destroy() NOEXCEPT override { delete this; } - - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override { - mPluginNamespace = pluginNamespace; - } - - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - //todo assert; - } - - IPluginV2 *clone() const NOEXCEPT override{ - ActivationLogisticRT *p = new ActivationLogisticRT(); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - int size; - -private: - std::string mPluginNamespace; }; - -class ActivationLogisticRTPluginCreator : public IPluginCreator{ -public: - ActivationLogisticRTPluginCreator(){ - mPluginAttributes.clear(); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override { - ActivationLogisticRT *pluginObj = new ActivationLogisticRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } - - IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override { - ActivationLogisticRT *pluginObj = new ActivationLogisticRT(); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ActivationLogisticRT_tkDNN"; - } - -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ActivationLogisticRTPluginCreator); \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ActivationMishRT.h b/include/tkDNN/pluginsRT/ActivationMishRT.h index a34418b..5d660af 100644 --- a/include/tkDNN/pluginsRT/ActivationMishRT.h +++ b/include/tkDNN/pluginsRT/ActivationMishRT.h @@ -1,134 +1,61 @@ #include #include "../kernels.h" -class ActivationMishRT : public IPluginV2 { +class ActivationMishRT : public IPlugin { public: - ActivationMishRT() {} - - ~ActivationMishRT() {} - - ActivationMishRT(const void *data, size_t length) { - const char *buf = reinterpret_cast(data), *bufCheck = buf; - size = readBUF(buf); - assert(buf == bufCheck + length); - } + ActivationMishRT() { - int getNbOutputs() const NOEXCEPT override { return 1; } + } - Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override { return inputs[0]; } + ~ActivationMishRT(){ - void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, - PluginFormat format, int maxBatchSize) NOEXCEPT override { - assert(format == PluginFormat::kLINEAR); - size = 1; - for (int i = 0; i < outputDims[0].nbDims; i++) - size *= outputDims[0].d[i]; - } + } - int initialize() NOEXCEPT override { return 0; } + int getNbOutputs() const override { + return 1; + } - virtual void terminate() NOEXCEPT override {} + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return inputs[0]; + } - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0; } + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + size = 1; + for(int i=0; i(inputs[0]), - reinterpret_cast(outputs[0]), batchSize * size, stream); - return 0; - } + 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 { + + activationMishForward((dnnType*)reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize*size, stream); + return 0; + } - virtual size_t getSerializationSize() const NOEXCEPT override { - return 1 * sizeof(int); - } + virtual size_t getSerializationSize() override { + return 1*sizeof(int); + } - virtual void serialize(void *buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer), *a = buf; - tk::dnn::writeBUF(buf, size); - assert(buf == a + getSerializationSize()); - } + virtual void serialize(void* buffer) override { + char *buf = reinterpret_cast(buffer),*a=buf; + tk::dnn::writeBUF(buf, size); + assert(buf == a + getSerializationSize()); + } - const char *getPluginType() const NOEXCEPT override { - return "ActivationMishRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override { - return "1"; - } - - void destroy() NOEXCEPT override { delete this; } - - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - } - - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char *plguinNamespace) NOEXCEPT override { - mPluginNamespace = plguinNamespace; - } - - IPluginV2 *clone() const NOEXCEPT override { - ActivationMishRT *p = new ActivationMishRT(); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - - int size; -private: - std::string mPluginNamespace; + int size; }; - -class ActivationMishRTPluginCreator : public IPluginCreator { -public: - ActivationMishRTPluginCreator() { - mPluginAttributes.clear(); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override { - mPluginNamespace = pluginNamespace; - } - - const char* getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char* name,const void* serialData,size_t serialLength) NOEXCEPT override{ - ActivationMishRT *pluginObj = new ActivationMishRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - ActivationMishRT *pluginObj = new ActivationMishRT(); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ActivationMishRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ActivationMishRTPluginCreator); \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h index 4d2652e..50ceb81 100644 --- a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h +++ b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h @@ -1,149 +1,63 @@ #include #include "../kernels.h" - -class ActivationReLUCeiling : public IPluginV2 { +class ActivationReLUCeiling : public IPlugin { public: - ActivationReLUCeiling(const float ceiling) { - this->ceiling = ceiling; - } + ActivationReLUCeiling(const float ceiling) { + this->ceiling = ceiling; + } - ~ActivationReLUCeiling() { + ~ActivationReLUCeiling(){ - } + } - ActivationReLUCeiling(const void *data, size_t length) { - const char *buf = reinterpret_cast(data), *bufCheck = buf; - ceiling = readBUF(buf); - size = readBUF(buf); - assert(buf == bufCheck + length); - } + int getNbOutputs() const override { + return 1; + } - int getNbOutputs() const NOEXCEPT override { - return 1; - } + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return inputs[0]; + } - Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override { - return inputs[0]; - } + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + size = 1; + for(int i=0; i(inputs[0]), - reinterpret_cast(outputs[0]), batchSize * size, ceiling, stream); - return 0; - } + virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { + + activationReLUCeilingForward((dnnType*)reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize*size, ceiling, stream); + return 0; + } - virtual size_t getSerializationSize() const NOEXCEPT override { - return 1 * sizeof(int) + 1 * sizeof(float); - } + virtual size_t getSerializationSize() override { + return 1*sizeof(int) + 1*sizeof(float); + } - virtual void serialize(void *buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer), *a = buf; - tk::dnn::writeBUF(buf, ceiling); - tk::dnn::writeBUF(buf, size); - assert(buf = a + getSerializationSize()); + virtual void serialize(void* buffer) override { + char *buf = reinterpret_cast(buffer),*a=buf; + tk::dnn::writeBUF(buf, ceiling); + tk::dnn::writeBUF(buf, size); + assert(buf = a + getSerializationSize()); + + } - } - - IPluginV2 *clone() const NOEXCEPT override { - ActivationReLUCeiling *p = new ActivationReLUCeiling(ceiling); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - - bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override { - return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR); - } - - void destroy() NOEXCEPT override { delete this; }; - - const char *getPluginType() const NOEXCEPT override { - return "ActivationReLUCeilingRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override { - return "1"; - } - - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override { - mPluginNamespace = pluginNamespace; - } - - int size; - float ceiling; -private: - std::string mPluginNamespace; + int size; + float ceiling; }; - -class ActivationReLUCeilingPluginCreator : public IPluginCreator { -public: - ActivationReLUCeilingPluginCreator() { - mPluginAttributes.emplace_back(PluginField("ceiling", nullptr, PluginFieldType::kFLOAT32, 1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override { - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override { - ActivationReLUCeiling *pluginObj = new ActivationReLUCeiling(serialData, serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char *name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - float ceiling = *(static_cast(fields[0].data)); - ActivationReLUCeiling *pluginObj = new ActivationReLUCeiling(ceiling); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ActivationReLUCeilingRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ActivationReLUCeilingPluginCreator); diff --git a/include/tkDNN/pluginsRT/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h index 46aa9e7..711198f 100644 --- a/include/tkDNN/pluginsRT/DeformableConvRT.h +++ b/include/tkDNN/pluginsRT/DeformableConvRT.h @@ -1,7 +1,8 @@ #include #include "../kernels.h" -class DeformableConvRT : public IPluginV2 { + +class DeformableConvRT : public IPlugin { @@ -26,8 +27,6 @@ public: this->o_c = o_c; this->o_h = o_h; this->o_w = o_w; - this->defRT = deformable; - height_ones = (i_h + 2 * ph - (1 * (kh - 1) + 1)) / sh + 1; width_ones = (i_w + 2 * pw - (1 * (kw - 1) + 1)) / sw + 1; dim_ones = i_c * kh * kw * 1 * height_ones * width_ones; @@ -39,6 +38,7 @@ public: checkCuda( cudaMalloc(&mask, chunk_dim*sizeof(dnnType))); checkCuda( cudaMalloc(&ones_d2, dim_ones*sizeof(dnnType))); if(deformable != nullptr) { + this->defRT = deformable; checkCuda( cudaMemcpy(data_d, deformable->data_d, sizeof(dnnType)*i_c * o_c * kh * kw * 1, cudaMemcpyDeviceToDevice) ); checkCuda( cudaMemcpy(bias2_d, deformable->bias2_d, sizeof(dnnType)*o_c, cudaMemcpyDeviceToDevice) ); checkCuda( cudaMemcpy(ones_d1, deformable->ones_d1, sizeof(dnnType)*height_ones*width_ones, cudaMemcpyDeviceToDevice) ); @@ -61,78 +61,27 @@ public: cublasDestroy(handle); } - DeformableConvRT(const void *data,size_t length){ - const char* buf = reinterpret_cast(data),*bufCheck = buf; - chunk_dim = readBUF(buf); - kh = readBUF(buf); - kw = readBUF(buf); - sh = readBUF(buf); - sw = readBUF(buf); - ph = readBUF(buf); - pw = readBUF(buf); - deformableGroup = readBUF(buf); - i_n = readBUF(buf); - i_c = readBUF(buf); - i_h = readBUF(buf); - i_w = readBUF(buf); - o_n = readBUF(buf); - o_c = readBUF(buf); - o_h = readBUF(buf); - o_w = readBUF(buf); - dnnType *aus = new dnnType[chunk_dim*2]; - for(int i=0;i(buf); - checkCuda(cudaMemcpy(offset,aus,sizeof(dnnType)*2*chunk_dim,cudaMemcpyHostToDevice)); - free(aus); - - aus = new dnnType[chunk_dim]; - for(int i=0;i(buf); - checkCuda(cudaMemcpy(mask,aus,sizeof(dnnType)*chunk_dim,cudaMemcpyHostToDevice)); - free(aus); - - aus = new dnnType[i_c*o_c*kh*kw*1]; - for(int i=0;i<(i_c*o_c*kh*kw*1);i++) - aus[i] = readBUF(buf); - checkCuda(cudaMemcpy(data_d,aus,sizeof(dnnType)*(i_c*o_c*kh*kw*1),cudaMemcpyHostToDevice)); - free(aus); - - aus = new dnnType[o_c]; - for(int i=0; i < o_c; i++) - aus[i] = readBUF(buf); - checkCuda( cudaMemcpy(bias2_d, aus, sizeof(dnnType)*o_c, cudaMemcpyHostToDevice) ); - free(aus); - - aus = new dnnType[height_ones * width_ones]; - for(int i=0; i(buf); - checkCuda( cudaMemcpy(ones_d1, aus, sizeof(dnnType)*height_ones * width_ones, cudaMemcpyHostToDevice) ); - free(aus); - - aus = new dnnType[dim_ones]; - for(int i=0; i(buf); - checkCuda( cudaMemcpy(ones_d2, aus, sizeof(dnnType)*dim_ones, cudaMemcpyHostToDevice) ); - free(aus); - - assert(buf == bufCheck + length); + int getNbOutputs() const override { + return 1; } - int getNbOutputs() const NOEXCEPT override {return 1;} - - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w}; + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return DimsCHW{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w}; } - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format,int maxBatchSize) NOEXCEPT override { } + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { } - int initialize() NOEXCEPT override {return 0;} + int initialize() override { + return 0; + } - virtual void terminate() NOEXCEPT override { } + virtual void terminate() override { } - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} + virtual size_t getWorkspaceSize(int maxBatchSize) const override { + return 0; + } - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { + virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); dnnType *output_conv = (dnnType*)reinterpret_cast(inputs[1]); @@ -160,12 +109,13 @@ public: return 0; } - virtual size_t getSerializationSize() const NOEXCEPT override { + + virtual size_t getSerializationSize() override { return 16 * sizeof(int) + chunk_dim * 3 * sizeof(dnnType) + (i_c * o_c * kh * kw * 1 ) * sizeof(dnnType) + o_c * sizeof(dnnType) + height_ones * width_ones * sizeof(dnnType) + dim_ones * sizeof(dnnType); } - virtual void serialize(void* buffer) const NOEXCEPT override { + virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, chunk_dim); tk::dnn::writeBUF(buf, kh); @@ -216,35 +166,6 @@ public: assert(buf == a + getSerializationSize()); } - void destroy() NOEXCEPT override {delete this;} - - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - //todo assert - } - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginType() const NOEXCEPT override{ - return "DeformableConvRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - IPluginV2* clone() const NOEXCEPT override{ - DeformableConvRT *p = new DeformableConvRT(chunk_dim,kh,kw,sh,sw,ph,pw,deformableGroup,i_n,i_c,i_h,i_w,o_n,o_c,o_h,o_w,defRT); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - - cublasStatus_t stat; cublasHandle_t handle; int i_n, i_c, i_h, i_w; @@ -272,90 +193,4 @@ public: tk::dnn::DeformConv2d *defRT; - -private: - std::string mPluginNamespace; }; - -class DeformableConvRTPluginCreator : public IPluginCreator{ -public: - DeformableConvRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("chunk_dim",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("kh",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("kw",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("sh",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("sw",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("ph",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("pw",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("deformableGroup",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("i_n",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("i_c",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("i_h",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("i_w",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_n",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_c",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_h",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_w",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("defRT",nullptr,PluginFieldType::kUNKNOWN,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char *name,const void *serialData,size_t serialLength) NOEXCEPT override{ - DeformableConvRT *pluginObj = new DeformableConvRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - int chunk_dim = *(static_cast(fields[0].data)); - int kh = *(static_cast(fields[1].data)); - int kw = *(static_cast(fields[2].data)); - int sh = *(static_cast(fields[3].data)); - int sw = *(static_cast(fields[4].data)); - int ph = *(static_cast(fields[5].data)); - int pw = *(static_cast(fields[6].data)); - int deformableGroup = *(static_cast(fields[7].data)); - int i_n = *(static_cast(fields[8].data)); - int i_c = *(static_cast(fields[9].data)); - int i_h = *(static_cast(fields[10].data)); - int i_w = *(static_cast(fields[11].data)); - int o_n = *(static_cast(fields[12].data)); - int o_c = *(static_cast(fields[13].data)); - int o_h = *(static_cast(fields[14].data)); - int o_w = *(static_cast(fields[14].data)); - DeformConv2d *defRT = const_cast(static_cast(fields[15].data)); - DeformableConvRT *pluginObj = new DeformableConvRT(chunk_dim,kh,kw,sh,sw,ph,pw,deformableGroup,i_n,i_c,i_h,i_w,o_n,o_c,o_h,o_w,defRT); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "DeformableConvRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(DeformableConvRTPluginCreator); - diff --git a/include/tkDNN/pluginsRT/FlattenConcatRT.h b/include/tkDNN/pluginsRT/FlattenConcatRT.h index c8b6c0f..51aa1ab 100644 --- a/include/tkDNN/pluginsRT/FlattenConcatRT.h +++ b/include/tkDNN/pluginsRT/FlattenConcatRT.h @@ -1,6 +1,6 @@ #include -class FlattenConcatRT : public IPluginV2 { +class FlattenConcatRT : public IPlugin { public: FlattenConcatRT() { @@ -11,29 +11,19 @@ public: } } - FlattenConcatRT(const void *data,size_t length){ - const char *buf = reinterpret_cast(data),*bufCheck=buf; - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - rows = readBUF(buf); - cols = readBUF(buf); - assert(buf == bufCheck + length); - } - ~FlattenConcatRT(){ } - int getNbOutputs() const NOEXCEPT override { + int getNbOutputs() const override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{ inputs[0].d[0] * inputs[0].d[1] * inputs[0].d[2], 1, 1}; + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return DimsCHW{ inputs[0].d[0] * inputs[0].d[1] * inputs[0].d[2], 1, 1}; } - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format,int maxBatchSize) NOEXCEPT override { + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { assert(nbOutputs == 1 && nbInputs ==1); rows = inputDims[0].d[0]; cols = inputDims[0].d[1] * inputDims[0].d[2]; @@ -42,13 +32,19 @@ public: w = 1; } - int initialize() NOEXCEPT override {return 0;} + int initialize() override { + return 0; + } - virtual void terminate() NOEXCEPT override { checkERROR(cublasDestroy(handle));} + virtual void terminate() override { + checkERROR(cublasDestroy(handle)); + } - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override {return 0;} + virtual size_t getWorkspaceSize(int maxBatchSize) const override { + return 0; + } - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { + virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); @@ -63,11 +59,12 @@ public: return 0; } - virtual size_t getSerializationSize() const NOEXCEPT override { + + virtual size_t getSerializationSize() override { return 5*sizeof(int); } - virtual void serialize(void* buffer) const NOEXCEPT override { + virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer),*a = buf; tk::dnn::writeBUF(buf, c); tk::dnn::writeBUF(buf, h); @@ -77,86 +74,8 @@ public: assert(buf == a + getSerializationSize()); } - void destroy() NOEXCEPT override{delete this;} - - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - } - - const char *getPluginType() const NOEXCEPT override{ - return "FlattenConcatRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - IPluginV2 *clone() const NOEXCEPT override { - FlattenConcatRT *p = new FlattenConcatRT(); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - int c, h, w; int rows, cols; cublasStatus_t stat; - cublasHandle_t handle; -private: - std::string mPluginNamespace; + cublasHandle_t handle; }; - -class FlattenConcatRTPluginCreator : public IPluginCreator{ -public: - FlattenConcatRTPluginCreator(){ - mPluginAttributes.clear(); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char *name,const void *serialData,size_t serialLength) NOEXCEPT override{ - FlattenConcatRT *pluginObj = new FlattenConcatRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char *name,const PluginFieldCollection *fc) NOEXCEPT override{ - FlattenConcatRT *pluginObj = new FlattenConcatRT(); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "FlattenConcatRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(FlattenConcatRTPluginCreator); \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h index 53fdd4e..0899a34 100644 --- a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h +++ b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h @@ -1,8 +1,7 @@ #include #include "../kernels.h" - -class MaxPoolFixedSizeRT : public IPluginV2 { +class MaxPoolFixedSizeRT : public IPlugin { public: MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding) { @@ -16,40 +15,32 @@ public: this->padding = padding; } - MaxPoolFixedSizeRT(const void *data,size_t length){ - const char *buf = reinterpret_cast(data),*bufCheck = buf; - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - n = readBUF(buf); - stride_H = readBUF(buf); - stride_W = readBUF(buf); - winSize = readBUF(buf); - padding = readBUF(buf); - assert(buf == bufCheck + length); - } - ~MaxPoolFixedSizeRT(){ } - int getNbOutputs() const NOEXCEPT override { + int getNbOutputs() const override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{this->c, this->h, this->w}; + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return DimsCHW{this->c, this->h, this->w}; } - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format,int maxBatchSize) NOEXCEPT override { + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { } - int initialize() NOEXCEPT override {return 0;} + int initialize() override { + return 0; + } - virtual void terminate() NOEXCEPT override {} + virtual void terminate() override { + } - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} + virtual size_t getWorkspaceSize(int maxBatchSize) const override { + return 0; + } - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { + 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]); @@ -59,11 +50,11 @@ public: } - virtual size_t getSerializationSize() const NOEXCEPT override { + virtual size_t getSerializationSize() override { return 8*sizeof(int); } - virtual void serialize(void* buffer) const NOEXCEPT override { + virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, this->c); @@ -77,106 +68,8 @@ public: assert(buf == a + getSerializationSize()); } - void destroy() NOEXCEPT override{delete this;} - - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - //todo assert - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - const char *getPluginType() const NOEXCEPT override{ - return "MaxPoolingFixedSizeRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - IPluginV2 *clone() const NOEXCEPT override{ - MaxPoolFixedSizeRT *p = new MaxPoolFixedSizeRT(c,h,w,n,stride_H,stride_W,winSize,padding); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - - int n, c, h, w; int stride_H, stride_W; int winSize; int padding; - -private: - std::string mPluginNamespace; }; - -class MaxPoolFixedSizeRTPluginCreator : public IPluginCreator{ -public: - MaxPoolFixedSizeRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("c",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("h",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("w",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("n",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("stride_H",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("stride_W",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("winSize",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("padding",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char *name,const void *serialData,size_t serialLength) NOEXCEPT override{ - MaxPoolFixedSizeRT *pluginObj = new MaxPoolFixedSizeRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char *name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - //todo assert - int c = *(static_cast(fields[0].data)); - int h = *(static_cast(fields[1].data)); - int w = *(static_cast(fields[2].data)); - int n = *(static_cast(fields[3].data)); - int stride_H = *(static_cast(fields[4].data)); - int stride_W = *(static_cast(fields[5].data)); - int winSize = *(static_cast(fields[6].data)); - int padding = *(static_cast(fields[7].data)); - MaxPoolFixedSizeRT *pluginObj = new MaxPoolFixedSizeRT(c,h,w,n,stride_H,stride_W,winSize,padding); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "MaxPoolingFixedSizeRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; - -}; - -REGISTER_TENSORRT_PLUGIN(MaxPoolFixedSizeRTPluginCreator); diff --git a/include/tkDNN/pluginsRT/RegionRT.h b/include/tkDNN/pluginsRT/RegionRT.h index 8e1c2d2..8487652 100644 --- a/include/tkDNN/pluginsRT/RegionRT.h +++ b/include/tkDNN/pluginsRT/RegionRT.h @@ -1,10 +1,11 @@ #include #include "../kernels.h" -class RegionRT : public IPluginV2 { +class RegionRT : public IPlugin { public: RegionRT(int classes, int coords, int num) { + this->classes = classes; this->coords = coords; this->num = num; @@ -14,39 +15,33 @@ public: } - RegionRT(const void *data,size_t length){ - const char *buf = reinterpret_cast(data),*bufCheck=buf; - classes = readBUF(buf); - coords = readBUF(buf); - num = readBUF(buf); - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - assert(buf == bufCheck+length); - } - - int getNbOutputs() const NOEXCEPT override { + int getNbOutputs() const override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { return inputs[0]; } - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format, int maxBatchSize) NOEXCEPT override { + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { c = inputDims[0].d[0]; h = inputDims[0].d[1]; w = inputDims[0].d[2]; } - int initialize() NOEXCEPT override { return 0; } + int initialize() override { + return 0; + } - virtual void terminate() NOEXCEPT override { } + virtual void terminate() override { + } - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0; } + virtual size_t getWorkspaceSize(int maxBatchSize) const override { + return 0; + } - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { + virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); @@ -73,11 +68,11 @@ public: } - virtual size_t getSerializationSize() const NOEXCEPT override { + virtual size_t getSerializationSize() override { return 6*sizeof(int); } - virtual void serialize(void* buffer) const NOEXCEPT override { + virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, classes); tk::dnn::writeBUF(buf, coords); @@ -88,34 +83,6 @@ public: assert(buf == a + getSerializationSize()); } - const char *getPluginType() const NOEXCEPT override{ - return "RegionRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - void destroy() NOEXCEPT override {delete this;} - - const char* getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - } - - IPluginV2* clone() const NOEXCEPT override{ - RegionRT *p = new RegionRT(classes,coords,num); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - int c, h, w; int classes, coords, num; @@ -125,64 +92,4 @@ public: return batch*c*h*w + n*w*h*(coords+classes+1) + entry*w*h + loc; } -private: - std::string mPluginNamespace; }; - -class RegionRTPluginCreator : public IPluginCreator{ -public: - RegionRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("classes",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("coords",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("num",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - IPluginV2 *deserializePlugin(const char* name,const void *serialData,size_t serialLength) NOEXCEPT override{ - RegionRT *pluginObj = new RegionRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - assert(fc->nbFields == 3); - assert(fields[0].type == PluginFieldType::kINT32); - assert(fields[1].type == PluginFieldType::kINT32); - assert(fields[2].type == PluginFieldType::kINT32); - int classes = *(static_cast(fields[0].data)); - int coords = *(static_cast(fields[1].data)); - int num = *(static_cast(fields[2].data)); - RegionRT *pluginObj = new RegionRT(classes,coords,num); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - - } - - const char *getPluginName() const NOEXCEPT override{ - return "RegionRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(RegionRTPluginCreator); - diff --git a/include/tkDNN/pluginsRT/ReorgRT.h b/include/tkDNN/pluginsRT/ReorgRT.h index 01e5468..c1b529a 100644 --- a/include/tkDNN/pluginsRT/ReorgRT.h +++ b/include/tkDNN/pluginsRT/ReorgRT.h @@ -1,7 +1,7 @@ #include #include "../kernels.h" -class ReorgRT : public IPluginV2 { +class ReorgRT : public IPlugin { public: ReorgRT(int stride) { @@ -12,34 +12,33 @@ public: } - ReorgRT(const void* data,size_t length){ - const char* buf = reinterpret_cast(data),*bufCheck = buf; - stride = readBUF(buf); - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - assert(buf == bufCheck + length); + int getNbOutputs() const override { + return 1; } - int getNbOutputs() const NOEXCEPT override {return 1;} - - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{inputs[0].d[0]*stride*stride, inputs[0].d[1]/stride, inputs[0].d[2]/stride}; + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return DimsCHW{inputs[0].d[0]*stride*stride, inputs[0].d[1]/stride, inputs[0].d[2]/stride}; } - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format, int maxBatchSize) NOEXCEPT override { + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { c = inputDims[0].d[0]; h = inputDims[0].d[1]; w = inputDims[0].d[2]; } - int initialize() NOEXCEPT override { return 0;} + int initialize() override { - virtual void terminate() NOEXCEPT override {} + return 0; + } - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} + virtual void terminate() override { + } - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT 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 { reorgForward((dnnType*)reinterpret_cast(inputs[0]), reinterpret_cast(outputs[0]), @@ -48,11 +47,11 @@ public: } - virtual size_t getSerializationSize() const NOEXCEPT override { + virtual size_t getSerializationSize() override { return 4*sizeof(int); } - virtual void serialize(void* buffer) const NOEXCEPT override { + virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, stride); tk::dnn::writeBUF(buf, c); @@ -60,84 +59,6 @@ public: tk::dnn::writeBUF(buf, w); assert(buf == a + getSerializationSize()); } - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{return true;} - - const char *getPluginType() const NOEXCEPT override{ - return "ReorgRT_tkDNN"; - } - - const char* getPluginVersion() const NOEXCEPT override{ - return "1"; - } - void destroy() NOEXCEPT override{ delete this;} - - const char* getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - IPluginV2* clone() const NOEXCEPT override{ - ReorgRT *p = new ReorgRT(stride); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } int c, h, w, stride; -private: - std::string mPluginNamespace; }; - -class ReorgRTPluginCreator : public IPluginCreator{ -public: - ReorgRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("stride",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char* getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2* deserializePlugin(const char* name,const void* serialData,size_t serialLength) NOEXCEPT override{ - ReorgRT *pluginObj = new ReorgRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection* fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - assert(fc->nbFields == 1); - assert(fields[0].type == PluginFieldType::kINT32); - int stride = *(static_cast(fields[0].data)); - ReorgRT *pluginObj = new ReorgRT(stride); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ReorgRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ReorgRTPluginCreator); - diff --git a/include/tkDNN/pluginsRT/ReshapeRT.h b/include/tkDNN/pluginsRT/ReshapeRT.h index fd74c61..37017c7 100644 --- a/include/tkDNN/pluginsRT/ReshapeRT.h +++ b/include/tkDNN/pluginsRT/ReshapeRT.h @@ -1,47 +1,42 @@ #include -class ReshapeRT : public IPluginV2 { +class ReshapeRT : public IPlugin { public: - ReshapeRT(dataDim_t newDim) { - new_dim = newDim; + ReshapeRT(dataDim_t new_dim) { n = new_dim.n; c = new_dim.c; h = new_dim.h; w = new_dim.w; } - ReshapeRT(const void *data,size_t length){ - const char *buf = reinterpret_cast(data),*bufCheck = buf; - new_dim.n = readBUF(buf); - new_dim.c = readBUF(buf); - new_dim.h = readBUF(buf); - new_dim.w = readBUF(buf); - assert(buf == bufCheck + length); - } - ~ReshapeRT(){ } - int getNbOutputs() const NOEXCEPT override { + int getNbOutputs() const override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{ c,h,w}; + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return DimsCHW{ c,h,w}; } - void configureWithFormat (const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, DataType type,PluginFormat format, int maxBatchSize) NOEXCEPT override { + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { } - int initialize() NOEXCEPT override {return 0;} + int initialize() override { + return 0; + } - virtual void terminate() NOEXCEPT override {} + virtual void terminate() override { + } - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} + virtual size_t getWorkspaceSize(int maxBatchSize) const override { + return 0; + } - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { + virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); @@ -49,11 +44,12 @@ public: return 0; } - virtual size_t getSerializationSize() const NOEXCEPT override { + + virtual size_t getSerializationSize() override { return 4*sizeof(int); } - virtual void serialize(void* buffer) const NOEXCEPT override { + virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer),*a = buf; tk::dnn::writeBUF(buf, n); tk::dnn::writeBUF(buf, c); @@ -62,87 +58,5 @@ public: assert(buf == a + getSerializationSize()); } - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - //todo assert - } - - const char *getPluginType() const NOEXCEPT override{ - return "1"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "ReshapeRT_tkDNN"; - } - - void destroy() NOEXCEPT override{delete this;} - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - IPluginV2 *clone() const NOEXCEPT override{ - ReshapeRT *p = new ReshapeRT(new_dim); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - int n, c, h, w; - dataDim_t new_dim; -private: - std::string mPluginNamespace; }; - -class ReshapeRTPluginCreator : public IPluginCreator{ -public: - ReshapeRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("new_dim",nullptr,PluginFieldType::kUNKNOWN,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char* name,const void *serialData,size_t serialLength) NOEXCEPT override{ - ReshapeRT *pluginObj = new ReshapeRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - dataDim_t newDim = *(static_cast(fields[0].data)); - ReshapeRT *pluginObj = new ReshapeRT(newDim); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ReshapeRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ReshapeRTPluginCreator); diff --git a/include/tkDNN/pluginsRT/ResizeLayerRT.h b/include/tkDNN/pluginsRT/ResizeLayerRT.h index 9a982fd..cde52bf 100644 --- a/include/tkDNN/pluginsRT/ResizeLayerRT.h +++ b/include/tkDNN/pluginsRT/ResizeLayerRT.h @@ -1,7 +1,7 @@ #include #include "../kernels.h" -class ResizeLayerRT : public IPluginV2 { +class ResizeLayerRT : public IPlugin { public: ResizeLayerRT(int c, int h, int w) { @@ -10,41 +10,35 @@ public: o_w = w; } - ResizeLayerRT(const void *data,size_t length){ - const char *buf = reinterpret_cast(data),*bufCheck = buf; - o_c = readBUF(buf); - o_h = readBUF(buf); - o_w = readBUF(buf); - i_c = readBUF(buf); - i_h = readBUF(buf); - i_w = readBUF(buf); - assert(buf == bufCheck + length); - } - ~ResizeLayerRT(){ } - int getNbOutputs() const NOEXCEPT override { + int getNbOutputs() const override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{o_c, o_h, o_w}; + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return DimsCHW{o_c, o_h, o_w}; } - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format,int maxBatchSize) NOEXCEPT override { + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { i_c = inputDims[0].d[0]; i_h = inputDims[0].d[1]; i_w = inputDims[0].d[2]; } - int initialize() NOEXCEPT override {return 0;} + int initialize() override { + return 0; + } - virtual void terminate() NOEXCEPT override {} + virtual void terminate() override { + } - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} + virtual size_t getWorkspaceSize(int maxBatchSize) const override { + return 0; + } - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { + virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { // printf("%d %d %d %d %d %d\n", i_c, i_w, i_h, o_c, o_w, o_h); resizeForward((dnnType*)reinterpret_cast(inputs[0]), reinterpret_cast(outputs[0]), @@ -53,11 +47,11 @@ public: } - virtual size_t getSerializationSize() const NOEXCEPT override { + virtual size_t getSerializationSize() override { return 6*sizeof(int); } - virtual void serialize(void* buffer) const NOEXCEPT override { + virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, o_c); @@ -70,96 +64,5 @@ public: assert(buf == a + getSerializationSize()); } - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - //todo assert - } - - const char *getPluginType() const NOEXCEPT override{ - return "ResizeLayerRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - void destroy() NOEXCEPT override{delete this;} - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - IPluginV2 *clone() const NOEXCEPT override{ - ResizeLayerRT *p = new ResizeLayerRT(o_c,o_h,o_w); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - int i_c, i_h, i_w, o_c, o_h, o_w; - -private: - std::string mPluginNamespace; }; - -class ResizeLayerRTPluginCreator : public IPluginCreator{ -public: - ResizeLayerRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("o_c",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_h",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_w",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char *name,const void *serialData,size_t serialLength) NOEXCEPT override{ - ResizeLayerRT *pluginObj = new ResizeLayerRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char *name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - assert(fc->nbFields == 3); - assert(fields[0].type == PluginFieldType::kINT32); - assert(fields[1].type == PluginFieldType::kINT32); - assert(fields[2].type == PluginFieldType::kINT32); - int oc = *(static_cast(fields[0].data)); - int oh = *(static_cast(fields[1].data)); - int ow = *(static_cast(fields[2].data)); - ResizeLayerRT *pluginObj = new ResizeLayerRT(oc,oh,ow); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ResizeLayerRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - - -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; - -}; - -REGISTER_TENSORRT_PLUGIN(ResizeLayerRTPluginCreator); - diff --git a/include/tkDNN/pluginsRT/RouteRT.h b/include/tkDNN/pluginsRT/RouteRT.h index 635566c..5a8c170 100644 --- a/include/tkDNN/pluginsRT/RouteRT.h +++ b/include/tkDNN/pluginsRT/RouteRT.h @@ -1,7 +1,7 @@ #include #include "../kernels.h" -class RouteRT : public IPluginV2 { +class RouteRT : public IPlugin { /** THIS IS NOT USED ANYMORE @@ -17,31 +17,17 @@ public: } - RouteRT(const void* data,size_t length){ - const char* buf = reinterpret_cast(data),*bufCheck = buf; - groups = readBUF(buf); - group_id = readBUF(buf); - in = readBUF(buf); - for(int i=0;i (buf); - } - c= readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - assert(buf == bufCheck + length); - } - - int getNbOutputs() const NOEXCEPT override { + int getNbOutputs() const override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { int out_c = 0; for(int i=0; i(outputs[0]); + for(int b=0; b(buffer),*a=buf; tk::dnn::writeBUF(buf, groups); tk::dnn::writeBUF(buf, group_id); @@ -92,90 +88,9 @@ public: assert(buf == a + getSerializationSize()); } - const char *getPluginType() const NOEXCEPT override{ - return "RouteRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - void destroy() NOEXCEPT override {delete this; } - - const char* getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override { return true;} - - IPluginV2* clone() const NOEXCEPT override{ - RouteRT *p = new RouteRT(groups,group_id); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } static const int MAX_INPUTS = 4; int in; int c_in[MAX_INPUTS]; int c, h, w; int groups, group_id; -private: - std::string mPluginNamespace; }; - -class RouteRTPluginCreator : public IPluginCreator{ -public: - RouteRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("groups",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("group_id",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char* name,const void* serialData,size_t serialLength) NOEXCEPT override{ - RouteRT *pluginObj = new RouteRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - assert(fc->nbFields == 2); - assert(fields[0].type == PluginFieldType::kINT32); - assert(fields[1].type == PluginFieldType::kINT32); - int groups = *(static_cast(fields[0].data)); - int group_id = *(static_cast(fields[1].data)); - RouteRT *pluginObj = new RouteRT(groups,group_id); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "RouteRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(RouteRTPluginCreator); diff --git a/include/tkDNN/pluginsRT/ShortcutRT.h b/include/tkDNN/pluginsRT/ShortcutRT.h index f5b8219..04091ac 100644 --- a/include/tkDNN/pluginsRT/ShortcutRT.h +++ b/include/tkDNN/pluginsRT/ShortcutRT.h @@ -1,52 +1,47 @@ #include #include "../kernels.h" - -class ShortcutRT : public IPluginV2 { +class ShortcutRT : public IPlugin { public: ShortcutRT(tk::dnn::dataDim_t bdim, bool mul) { - bDim = bdim; - this->bc = bDim.c; - this->bh = bDim.h; - this->bw = bDim.w; + this->bc = bdim.c; + this->bh = bdim.h; + this->bw = bdim.w; this->mul = mul; } - ~ShortcutRT(){} + ~ShortcutRT(){ - ShortcutRT(const void* data,size_t length){ - const char* buf =reinterpret_cast(data),*bufCheck = buf; - bDim.c = readBUF(buf); - bDim.h = readBUF(buf); - bDim.w = readBUF(buf); - bDim.l = 1; - mul = readBUF(buf); - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - assert(buf == bufCheck + length); } - int getNbOutputs() const NOEXCEPT override {return 1;} - - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]}; + int getNbOutputs() const override { + return 1; } - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format,int maxBatchSize) NOEXCEPT override { + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return DimsCHW{inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]}; + } + + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { c = inputDims[0].d[0]; h = inputDims[0].d[1]; w = inputDims[0].d[2]; } - int initialize() NOEXCEPT override {return 0;} + int initialize() override { - virtual void terminate() NOEXCEPT override {} + return 0; + } - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} + virtual void terminate() override { + } - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { + virtual size_t getWorkspaceSize(int maxBatchSize) const override { + return 0; + } + + virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); dnnType *srcDataBack = (dnnType*)reinterpret_cast(inputs[1]); @@ -59,11 +54,11 @@ public: } - virtual size_t getSerializationSize() const NOEXCEPT override { + virtual size_t getSerializationSize() override { return 6*sizeof(int) + sizeof(bool); } - virtual void serialize(void* buffer) const NOEXCEPT override { + virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, bc); tk::dnn::writeBUF(buf, bh); @@ -76,91 +71,7 @@ public: } - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - } - - const char* getPluginType() const NOEXCEPT override{ - return "1"; - } - - const char* getPluginVersion() const NOEXCEPT override{ - return "ShortcutRT_tkDNN"; - } - - void destroy() NOEXCEPT override{delete this;} - - const char* getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - IPluginV2 *clone() const NOEXCEPT override{ - ShortcutRT *p = new ShortcutRT(bDim,mul); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - int c, h, w; int bc, bh, bw; bool mul; - tk::dnn::dataDim_t bDim; -private: - std::string mPluginNamespace; }; - - -class ShortcutRTPluginCreator : public IPluginCreator { -public: - ShortcutRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("bDim",nullptr,PluginFieldType::kUNKNOWN,1)); - mPluginAttributes.emplace_back(PluginField("mul",nullptr,PluginFieldType::kUNKNOWN,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char *name,const void *serialData,size_t serialLength) NOEXCEPT override{ - ShortcutRT *pluginObj = new ShortcutRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char *name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - //todo assert - tk::dnn::dataDim_t bdim = *(static_cast(fields[0].data)); - bool mul = *(static_cast(fields[1].data)); - ShortcutRT *pluginObj = new ShortcutRT(bdim,mul); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ShortcutRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ShortcutRTPluginCreator); \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/UpsampleRT.h b/include/tkDNN/pluginsRT/UpsampleRT.h index 0a71c45..a11d7b4 100644 --- a/include/tkDNN/pluginsRT/UpsampleRT.h +++ b/include/tkDNN/pluginsRT/UpsampleRT.h @@ -1,47 +1,44 @@ #include #include "../kernels.h" - -class UpsampleRT : public IPluginV2 { +class UpsampleRT : public IPlugin { public: UpsampleRT(int stride) { this->stride = stride; } - UpsampleRT(const void *data,size_t length){ - const char* buf = reinterpret_cast(data),*bufCheck=buf; - stride = readBUF(buf); - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - assert(buf == bufCheck + length); + ~UpsampleRT(){ + } - - ~UpsampleRT(){} - - int getNbOutputs() const NOEXCEPT override { + int getNbOutputs() const override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3(inputs[0].d[0], inputs[0].d[1]*stride, inputs[0].d[2]*stride); + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return DimsCHW(inputs[0].d[0], inputs[0].d[1]*stride, inputs[0].d[2]*stride); } - void configureWithFormat (const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format,int maxBatchSize) NOEXCEPT override { + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { c = inputDims[0].d[0]; h = inputDims[0].d[1]; w = inputDims[0].d[2]; } - int initialize() NOEXCEPT override {return 0;} + int initialize() override { - virtual void terminate() NOEXCEPT override {} + return 0; + } - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} + virtual void terminate() override { + } - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { + virtual size_t getWorkspaceSize(int maxBatchSize) const override { + return 0; + } + + virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); @@ -52,9 +49,11 @@ public: } - virtual size_t getSerializationSize() const NOEXCEPT override { return 4*sizeof(int);} + virtual size_t getSerializationSize() override { + return 4*sizeof(int); + } - virtual void serialize(void* buffer) const NOEXCEPT override { + virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, stride); tk::dnn::writeBUF(buf, c); @@ -63,85 +62,5 @@ public: assert(buf == a + getSerializationSize()); } - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - //todo assert - return true; - } - - const char *getPluginType() const NOEXCEPT override{ - return "1"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "UpsampleRT_tkDNN"; - } - - void destroy() NOEXCEPT override{delete this;} - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - IPluginV2* clone() const NOEXCEPT override{ - UpsampleRT *p = new UpsampleRT(stride); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - int c, h, w, stride; -private: - std::string mPluginNamespace; }; - -class UpsampleRTPluginCreator : public IPluginCreator{ -public: - UpsampleRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("stride",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char* name,const void* serialData,size_t serialLength) NOEXCEPT override{ - UpsampleRT *pluginObj = new UpsampleRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - int stride = *(static_cast(fields[0].data)); - UpsampleRT *pluginObj = new UpsampleRT(stride); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "UpsampleRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(UpsampleRTPluginCreator); diff --git a/include/tkDNN/pluginsRT/YoloRT.h b/include/tkDNN/pluginsRT/YoloRT.h index cc866c6..5ffe39c 100644 --- a/include/tkDNN/pluginsRT/YoloRT.h +++ b/include/tkDNN/pluginsRT/YoloRT.h @@ -1,268 +1,143 @@ #include #include "../kernels.h" + #define YOLORT_CLASSNAME_W 256 -class YoloRT : public IPluginV2 { +class YoloRT : public IPlugin { + + public: - YoloRT(int classes, int num, tk::dnn::Yolo *Yolo = nullptr, int n_masks = 3, float scale_xy = 1, - float nms_thresh = 0.45, int nms_kind = 0, int new_coords = 0) { - this->yolo = Yolo; - this->classes = classes; - this->num = num; - this->n_masks = n_masks; - this->scaleXY = scale_xy; - this->nms_thresh = nms_thresh; - this->nms_kind = nms_kind; - this->new_coords = new_coords; + YoloRT(int classes, int num, tk::dnn::Yolo *yolo = nullptr, int n_masks=3, float scale_xy=1, float nms_thresh=0.45, int nms_kind=0, int new_coords=0) { + + this->classes = classes; + this->num = num; + this->n_masks = n_masks; + this->scaleXY = scale_xy; + this->nms_thresh = nms_thresh; + this->nms_kind = nms_kind; + this->new_coords = new_coords; mask = new dnnType[n_masks]; - bias = new dnnType[num * n_masks * 2]; - if (yolo != nullptr) { - memcpy(mask, yolo->mask_h, sizeof(dnnType) * n_masks); - memcpy(bias, yolo->bias_h, sizeof(dnnType) * num * n_masks * 2); - classesNames = yolo->classesNames; + bias = new dnnType[num*n_masks*2]; + if(yolo != nullptr) { + memcpy(mask, yolo->mask_h, sizeof(dnnType)*n_masks); + memcpy(bias, yolo->bias_h, sizeof(dnnType)*num*n_masks*2); + classesNames = yolo->classesNames; } - } + } - YoloRT(const void *data,size_t length){ - const char* buf = reinterpret_cast(data),*bufCheck = buf; - classes = readBUF(buf); - num = readBUF(buf); - n_masks = readBUF(buf); - scaleXY = readBUF(buf); - nms_thresh = readBUF(buf); - nms_kind = readBUF(buf); - new_coords = readBUF(buf); - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - for(int i=0;i(buf); - for(int i=0;i(buf); - classesNames.resize(classes); - for(int i=0;i(buf); - classesNames[1] = std::string(tmp); - } - assert(buf == bufCheck + length); - } + ~YoloRT(){ - ~YoloRT() { + } - } + int getNbOutputs() const override { + return 1; + } + + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return inputs[0]; + } + + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + c = inputDims[0].d[0]; + h = inputDims[0].d[1]; + w = inputDims[0].d[2]; + } + + int initialize() override { + + return 0; + } + + virtual void terminate() override { + } + + virtual size_t getWorkspaceSize(int maxBatchSize) const override { + return 0; + } + + virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { + + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + + checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); - int getNbOutputs() const NOEXCEPT override { - return 1; - } + for (int b = 0; b < batchSize; ++b){ + for(int n = 0; n < n_masks; ++n){ + int index = entry_index(b, n*w*h, 0); + if (new_coords == 1){ + if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); + } + else{ + activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); //x,y - Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override { - return inputs[0]; - } + if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); - void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, - PluginFormat format, int maxBatchSize) NOEXCEPT override { - c = inputDims[0].d[0]; - h = inputDims[0].d[1]; - w = inputDims[0].d[2]; - } - - int initialize() NOEXCEPT override { - - return 0; - } - - virtual void terminate() NOEXCEPT override { - } - - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { - return 0; - } - - virtual int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, - cudaStream_t stream) NOEXCEPT override { - - dnnType *srcData = (dnnType *) reinterpret_cast(inputs[0]); - dnnType *dstData = reinterpret_cast(outputs[0]); - - checkCuda(cudaMemcpyAsync(dstData, srcData, batchSize * c * h * w * sizeof(dnnType), cudaMemcpyDeviceToDevice, - stream)); - - - for (int b = 0; b < batchSize; ++b) { - for (int n = 0; n < n_masks; ++n) { - int index = entry_index(b, n * w * h, 0); - if (new_coords == 1) { - if (this->scaleXY != 1) - scalAdd(dstData + index, 2 * w * h, this->scaleXY, -0.5 * (this->scaleXY - 1), 1); - } else { - activationLOGISTICForward(srcData + index, dstData + index, 2 * w * h, stream); //x,y - - if (this->scaleXY != 1) - scalAdd(dstData + index, 2 * w * h, this->scaleXY, -0.5 * (this->scaleXY - 1), 1); - - index = entry_index(b, n * w * h, 4); - activationLOGISTICForward(srcData + index, dstData + index, (1 + classes) * w * h, stream); + index = entry_index(b, n*w*h, 4); + activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream); } } } - //std::cout<<"YOLO END\n"; - return 0; - } + //std::cout<<"YOLO END\n"; + return 0; + } - virtual size_t getSerializationSize() const NOEXCEPT override { - return 8 * sizeof(int) + 2 * sizeof(float) + n_masks * sizeof(dnnType) + num * n_masks * 2 * sizeof(dnnType) + - YOLORT_CLASSNAME_W * classes * sizeof(char); - } + virtual size_t getSerializationSize() override { + return 8*sizeof(int) + 2*sizeof(float)+ n_masks*sizeof(dnnType) + num*n_masks*2*sizeof(dnnType) + YOLORT_CLASSNAME_W*classes*sizeof(char); + } - bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override { - return true; //todo implement proper supportsFormat - } + virtual void serialize(void* buffer) override { + char *buf = reinterpret_cast(buffer),*a=buf; + tk::dnn::writeBUF(buf, classes); //std::cout << "Classes :" << classes << std::endl; + tk::dnn::writeBUF(buf, num); //std::cout << "Num : " << num << std::endl; + tk::dnn::writeBUF(buf, n_masks); //std::cout << "N_Masks" << n_masks << std::endl; + tk::dnn::writeBUF(buf, scaleXY); //std::cout << "ScaleXY :" << scaleXY << std::endl; + tk::dnn::writeBUF(buf, nms_thresh); //std::cout << "nms_thresh :" << nms_thresh << std::endl; + tk::dnn::writeBUF(buf, nms_kind); //std::cout << "nms_kind : " << nms_kind << std::endl; + tk::dnn::writeBUF(buf, new_coords); //std::cout << "new_coords : " << new_coords << std::endl; + tk::dnn::writeBUF(buf, c); //std::cout << "C : " << c << std::endl; + tk::dnn::writeBUF(buf, h); //std::cout << "H : " << h << std::endl; + tk::dnn::writeBUF(buf, w); //std::cout << "C : " << c << std::endl; + for (int i = 0; i < n_masks; i++) + { + tk::dnn::writeBUF(buf, mask[i]); //std::cout << "mask[i] : " << mask[i] << std::endl; + } + for (int i = 0; i < n_masks * 2 * num; i++) + { + tk::dnn::writeBUF(buf, bias[i]); //std::cout << "bias[i] : " << bias[i] << std::endl; + } - virtual void serialize(void *buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer), *a = buf; - tk::dnn::writeBUF(buf, classes); //std::cout << "Classes :" << classes << std::endl; - tk::dnn::writeBUF(buf, num); //std::cout << "Num : " << num << std::endl; - tk::dnn::writeBUF(buf, n_masks); //std::cout << "N_Masks" << n_masks << std::endl; - tk::dnn::writeBUF(buf, scaleXY); //std::cout << "ScaleXY :" << scaleXY << std::endl; - tk::dnn::writeBUF(buf, nms_thresh); //std::cout << "nms_thresh :" << nms_thresh << std::endl; - tk::dnn::writeBUF(buf, nms_kind); //std::cout << "nms_kind : " << nms_kind << std::endl; - tk::dnn::writeBUF(buf, new_coords); //std::cout << "new_coords : " << new_coords << std::endl; - tk::dnn::writeBUF(buf, c); //std::cout << "C : " << c << std::endl; - tk::dnn::writeBUF(buf, h); //std::cout << "H : " << h << std::endl; - tk::dnn::writeBUF(buf, w); //std::cout << "C : " << c << std::endl; - for (int i = 0; i < n_masks; i++) { - tk::dnn::writeBUF(buf, mask[i]); //std::cout << "mask[i] : " << mask[i] << std::endl; - } - for (int i = 0; i < n_masks * 2 * num; i++) { - tk::dnn::writeBUF(buf, bias[i]); //std::cout << "bias[i] : " << bias[i] << std::endl; - } + // save classes names + for(int i=0; isetPluginNamespace(mPluginNamespace.c_str()); - return p; - } - - Yolo *yolo; - int c, h, w; + int c, h, w; int classes, num, n_masks; - float scaleXY; - float nms_thresh; - int nms_kind; - int new_coords; - std::vector classesNames; + float scaleXY; + float nms_thresh; + int nms_kind; + int new_coords; + std::vector classesNames; dnnType *mask; dnnType *bias; - int entry_index(int batch, int location, int entry) { - int n = location / (w * h); - int loc = location % (w * h); - return batch * c * h * w + n * w * h * (4 + classes + 1) + entry * w * h + loc; - } - -private: - std::string mPluginNamespace; + int entry_index(int batch, int location, int entry) { + int n = location / (w*h); + int loc = location % (w*h); + return batch*c*h*w + n*w*h*(4+classes+1) + entry*w*h + loc; + } }; - -class YoloRTPluginCreator : public IPluginCreator{ -public: - YoloRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("classes",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("num",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("yolo",nullptr,PluginFieldType::kUNKNOWN,1)); - mPluginAttributes.emplace_back(PluginField("numMasks",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("scaleXY",nullptr,PluginFieldType::kFLOAT32,1)); - mPluginAttributes.emplace_back(PluginField("nmsThresh",nullptr,PluginFieldType::kFLOAT32,1)); - mPluginAttributes.emplace_back(PluginField("nmsKind",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("newCoords",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char *name,const void *serialData,size_t serialLength) NOEXCEPT override{ - YoloRT *pluginObj = new YoloRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - //todo assert - int classes = *(static_cast(fields[0].data)); - int num = *(static_cast(fields[1].data)); - Yolo *yoloTemp = const_cast(static_cast(fields[2].data)); - int numMasks = *(static_cast(fields[3].data)); - float scaleXY = *(static_cast(fields[4].data)); - float nmsThresh = *(static_cast(fields[5].data)); - int nmsKind = *(static_cast(fields[6].data)); - int newCoords = *(static_cast(fields[7].data)); - YoloRT *pluginObj = new YoloRT(classes,num,yoloTemp,numMasks,scaleXY,nmsThresh,nmsKind,newCoords); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "YoloRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(YoloRTPluginCreator); diff --git a/include/tkDNN/utils.h b/include/tkDNN/utils.h index 8041f3b..65375ba 100644 --- a/include/tkDNN/utils.h +++ b/include/tkDNN/utils.h @@ -11,7 +11,6 @@ #include "cuda_runtime_api.h" #include #include -#include #ifdef __linux__ @@ -23,15 +22,6 @@ #include - - -#if NV_TENSORRT_MAJOR > 7 -#define NOEXCEPT noexcept -#else -#define NOEXCEPT -#endif - - #define dnnType float diff --git a/src/Int8Calibrator.cpp b/src/Int8Calibrator.cpp index 9fab18f..773a9d8 100644 --- a/src/Int8Calibrator.cpp +++ b/src/Int8Calibrator.cpp @@ -8,7 +8,7 @@ Int8EntropyCalibrator::Int8EntropyCalibrator(BatchStream& stream, int firstBatch mCalibTableFilePath(calibTableFilePath), mInputBlobName(inputBlobName.c_str()), mReadCache(readCache) { - nvinfer1::Dims4 dims = mStream.getDims(); + nvinfer1::DimsNCHW dims = mStream.getDims(); mInputCount = mStream.getBatchSize() * dims.c() * dims.h() * dims.w(); checkCuda(cudaMalloc(&mDeviceInput, mInputCount * sizeof(float))); mStream.reset(firstBatch); From 33bbcd216cd6cf72f3dc7f99db2207b56d733ba6 Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Thu, 4 Nov 2021 19:28:55 +0100 Subject: [PATCH 4/6] Fix max elem (remove thrust) for segmentation Signed-off-by: Micaela Verucchi --- src/kernels/postprocessing.cu | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/src/kernels/postprocessing.cu b/src/kernels/postprocessing.cu index 63643eb..c175e1f 100644 --- a/src/kernels/postprocessing.cu +++ b/src/kernels/postprocessing.cu @@ -46,11 +46,16 @@ void maxElem_kernel(float *src_begin, float *dst_begin, const int n_classes, con if (i > size) return; - thrust::device_ptr dPbeg ( &src_begin[i*n_classes] ) ; - thrust::device_ptr dPend = dPbeg + n_classes; - thrust::device_ptr result = thrust::max_element(thrust::device,dPbeg, dPend); + float max = 0; + int max_idx = 0; + for( int j = i*n_classes; j < i*n_classes + n_classes; ++j ){ + if( src_begin[j] > max ){ + max = src_begin[j]; + max_idx = j; + } + } - dst_begin[i] = result - dPbeg; + dst_begin[i] = max_idx - i*n_classes; } void maxElem(dnnType *src_begin, dnnType *dst_begin, const int c, const int h, const int w){ From 9ba6d842820ddd2b7e4220a9697121d2a5463137 Mon Sep 17 00:00:00 2001 From: Harshvardhan Chandirasekar <43143075+perseusdg@users.noreply.github.com> Date: Tue, 11 Jan 2022 13:26:53 +0530 Subject: [PATCH 5/6] Create codeql-analysis.yml --- codeql-analysis.yml | 70 +++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 70 insertions(+) create mode 100644 codeql-analysis.yml diff --git a/codeql-analysis.yml b/codeql-analysis.yml new file mode 100644 index 0000000..fae285c --- /dev/null +++ b/codeql-analysis.yml @@ -0,0 +1,70 @@ +# For most projects, this workflow file will not need changing; you simply need +# to commit it to your repository. +# +# You may wish to alter this file to override the set of languages analyzed, +# or to provide custom queries or build logic. +# +# ******** NOTE ******** +# We have attempted to detect the languages in your repository. Please check +# the `language` matrix defined below to confirm you have the correct set of +# supported CodeQL languages. +# +name: "CodeQL" + +on: + push: + branches: [ tensorrt8,depth,dev ] + pull_request: + # The branches below must be a subset of the branches above + branches: [ tensorrt8,depth,dev ] + schedule: + - cron: '44 13 * * 2' + +jobs: + analyze: + name: Analyze + runs-on: ubuntu-lts + permissions: + actions: read + contents: read + security-events: write + + strategy: + fail-fast: false + matrix: + language: [ 'cpp', 'python' ] + # CodeQL supports [ 'cpp', 'csharp', 'go', 'java', 'javascript', 'python', 'ruby' ] + # Learn more about CodeQL language support at https://git.io/codeql-language-support + + steps: + - name: Checkout repository + uses: actions/checkout@v2 + + # Initializes the CodeQL tools for scanning. + - name: Initialize CodeQL + uses: github/codeql-action/init@v1 + with: + languages: ${{ matrix.language }} + # If you wish to specify custom queries, you can do so here or in a config file. + # By default, queries listed here will override any specified in a config file. + # Prefix the list here with "+" to use these queries and those in the config file. + # queries: ./path/to/local/query, your-org/your-repo/queries@main + + # Autobuild attempts to build any compiled languages (C/C++, C#, or Java). + # If this step fails, then you should remove it and run the build manually (see below) + - name: Autobuild + uses: github/codeql-action/autobuild@v1 + + # ℹī¸ Command-line programs to run using the OS shell. + # 📚 https://git.io/JvXDl + + # ✏ī¸ If the Autobuild fails above, remove it and uncomment the following three lines + # and modify them (or add more) to build your code if your project + # uses a compiled language + + #- run: | + # make bootstrap + # make release + + - name: Perform CodeQL Analysis + uses: github/codeql-action/analyze@v1 From 2f80c3c38d29a07846f6b5c9f15c787b681ca9ef Mon Sep 17 00:00:00 2001 From: Harshvardhan Chandirasekar <43143075+perseusdg@users.noreply.github.com> Date: Tue, 11 Jan 2022 13:27:55 +0530 Subject: [PATCH 6/6] Delete codeql-analysis.yml --- codeql-analysis.yml | 70 --------------------------------------------- 1 file changed, 70 deletions(-) delete mode 100644 codeql-analysis.yml diff --git a/codeql-analysis.yml b/codeql-analysis.yml deleted file mode 100644 index fae285c..0000000 --- a/codeql-analysis.yml +++ /dev/null @@ -1,70 +0,0 @@ -# For most projects, this workflow file will not need changing; you simply need -# to commit it to your repository. -# -# You may wish to alter this file to override the set of languages analyzed, -# or to provide custom queries or build logic. -# -# ******** NOTE ******** -# We have attempted to detect the languages in your repository. Please check -# the `language` matrix defined below to confirm you have the correct set of -# supported CodeQL languages. -# -name: "CodeQL" - -on: - push: - branches: [ tensorrt8,depth,dev ] - pull_request: - # The branches below must be a subset of the branches above - branches: [ tensorrt8,depth,dev ] - schedule: - - cron: '44 13 * * 2' - -jobs: - analyze: - name: Analyze - runs-on: ubuntu-lts - permissions: - actions: read - contents: read - security-events: write - - strategy: - fail-fast: false - matrix: - language: [ 'cpp', 'python' ] - # CodeQL supports [ 'cpp', 'csharp', 'go', 'java', 'javascript', 'python', 'ruby' ] - # Learn more about CodeQL language support at https://git.io/codeql-language-support - - steps: - - name: Checkout repository - uses: actions/checkout@v2 - - # Initializes the CodeQL tools for scanning. - - name: Initialize CodeQL - uses: github/codeql-action/init@v1 - with: - languages: ${{ matrix.language }} - # If you wish to specify custom queries, you can do so here or in a config file. - # By default, queries listed here will override any specified in a config file. - # Prefix the list here with "+" to use these queries and those in the config file. - # queries: ./path/to/local/query, your-org/your-repo/queries@main - - # Autobuild attempts to build any compiled languages (C/C++, C#, or Java). - # If this step fails, then you should remove it and run the build manually (see below) - - name: Autobuild - uses: github/codeql-action/autobuild@v1 - - # ℹī¸ Command-line programs to run using the OS shell. - # 📚 https://git.io/JvXDl - - # ✏ī¸ If the Autobuild fails above, remove it and uncomment the following three lines - # and modify them (or add more) to build your code if your project - # uses a compiled language - - #- run: | - # make bootstrap - # make release - - - name: Perform CodeQL Analysis - uses: github/codeql-action/analyze@v1