From f40cee71cc60f2f80b2afb0c721ae1fc92c3eb72 Mon Sep 17 00:00:00 2001 From: Reese Wang Date: Fri, 18 Jun 2021 14:44:08 +0800 Subject: [PATCH 01/46] add trt LT version helper --- paddle/fluid/inference/tensorrt/helper.h | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/paddle/fluid/inference/tensorrt/helper.h b/paddle/fluid/inference/tensorrt/helper.h index 6158fd130bad8..8b557c6e6b8ce 100644 --- a/paddle/fluid/inference/tensorrt/helper.h +++ b/paddle/fluid/inference/tensorrt/helper.h @@ -31,6 +31,10 @@ namespace tensorrt { ((NV_TENSORRT_MAJOR * 1000 + NV_TENSORRT_MINOR * 100 + \ NV_TENSORRT_PATCH * 10 + NV_TENSORRT_BUILD) >= version) +#define IS_TRT_VERSION_LT(version) \ + ((NV_TENSORRT_MAJOR * 1000 + NV_TENSORRT_MINOR * 100 + \ + NV_TENSORRT_PATCH * 10 + NV_TENSORRT_BUILD) < version) + #define TRT_VERSION \ NV_TENSORRT_MAJOR * 1000 + NV_TENSORRT_MINOR * 100 + \ NV_TENSORRT_PATCH * 10 + NV_TENSORRT_BUILD From 75d20da9581c6ccb83c1a9ab5b4ba2b681cd57a1 Mon Sep 17 00:00:00 2001 From: Reese Wang Date: Sun, 20 Jun 2021 16:55:37 +0800 Subject: [PATCH 02/46] upgrade PluginTensorRT to IPluginV2Ext --- .../inference/tensorrt/plugin/trt_plugin.cc | 28 +++++- .../inference/tensorrt/plugin/trt_plugin.h | 85 ++++++++++++++----- 2 files changed, 89 insertions(+), 24 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc index 55bc786746bea..751deac45859b 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc @@ -49,7 +49,7 @@ inline size_t SeriaSize(const std::vector& input_dims, SerializedSize(with_fp16)); } -void PluginTensorRT::serializeBase(void*& buffer) { +void PluginTensorRT::serializeBase(void*& buffer) const { Seria(buffer, input_dims_, max_batch_size_, data_type_, data_format_, with_fp16_); } @@ -60,7 +60,7 @@ void PluginTensorRT::deserializeBase(void const*& serial_data, &data_type_, &data_format_, &with_fp16_); } -size_t PluginTensorRT::getBaseSerializationSize() { +size_t PluginTensorRT::getBaseSerializationSize() const { return SeriaSize(input_dims_, max_batch_size_, data_type_, data_format_, with_fp16_); } @@ -81,6 +81,30 @@ void PluginTensorRT::configureWithFormat( max_batch_size_ = max_batch_size; } +nvinfer1::DataType PluginTensorRT::getOutputDataType( + int32_t index, const nvinfer1::DataType* input_types, + int32_t nb_inputs) const { + return input_types[0]; +} + +bool PluginTensorRT::isOutputBroadcastAcrossBatch( + int32_t output_index, const bool* input_is_broadcasted, + int32_t nb_inputs) const { + return false; +} + +bool PluginTensorRT::canBroadcastInputAcrossBatch(int32_t input_index) const { + return false; +} + +void PluginTensorRT::configurePlugin( + const nvinfer1::Dims* input_dims, int32_t nb_inputs, + const nvinfer1::Dims* output_dims, int32_t nb_outputs, + const nvinfer1::DataType* input_types, + const nvinfer1::DataType* output_types, const bool* input_is_broadcast, + const bool* output_is_broadcast, nvinfer1::PluginFormat float_format, + int32_t max_batch_size) {} + void PluginTensorRTV2Ext::serializeBase(void*& buffer) const { Seria(buffer, input_dims_, max_batch_size_, data_type_, data_format_, with_fp16_); diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h index ce3133ae99e94..428e1dfc50eb8 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h @@ -45,71 +45,109 @@ typedef std::function typedef std::function PluginConstructFunc; // Deprecated. Do not inherit this class, please refer to PluginTensorRTV2Ext -class PluginTensorRT : public nvinfer1::IPluginExt { +class PluginTensorRT : public nvinfer1::IPluginV2Ext { public: PluginTensorRT() : with_fp16_(false) {} // It was used for TensorRT deserialization. // It should not be called by users. PluginTensorRT(const void* serialized_data, size_t length) {} + virtual ~PluginTensorRT() {} nvinfer1::Dims const& getInputDims(int index) const { return input_dims_.at(index); } + size_t getMaxBatchSize() const { return max_batch_size_; } + nvinfer1::DataType getDataType() const { return data_type_; } + nvinfer1::PluginFormat getDataFormat() const { return data_format_; } - virtual const char* getPluginVersion() const { return "1"; } void AddInput(nvinfer1::ITensor* input) { inputs_.push_back(input); } + std::vector& GetInputs() { return inputs_; } - virtual nvinfer1::IPluginExt* clone() const = 0; + // IPluginV2 virtual const char* getPluginType() const = 0; - // Following functions are inherit from nvinfer1::IPluginExt - // Get the number of outputs from the layer + virtual const char* getPluginVersion() const { return "1"; } + int getNbOutputs() const { return 1; } - // Get the dimension of an output tensor + virtual nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* input_dims, int num_inputs) = 0; - // Find the workspace size required by the layer - size_t getWorkspaceSize(int) const override { return 0; } + + // Check format support. The default is FLOAT32 and NCHW. + bool supportsFormat(nvinfer1::DataType type, + nvinfer1::PluginFormat format) const override; + + // Configure the layer + void configureWithFormat(const nvinfer1::Dims* input_dims, int num_inputs, + const nvinfer1::Dims* output_dims, int num_outputs, + nvinfer1::DataType type, + nvinfer1::PluginFormat format, + int max_batch_size) override; // Initialize the layer for execution. - // This is called when the engine is created. int initialize() override { return 0; } + // Shutdown the layer. This is called when the engine is destroyed void terminate() override {} + + // Find the workspace size required by the layer + size_t getWorkspaceSize(int) const override { return 0; } + // Execute the layer virtual int enqueue(int batch_size, const void* const* inputs, void** outputs, void* workspace, cudaStream_t stream) = 0; // Find the size of the serialization buffer required - virtual size_t getSerializationSize() = 0; + virtual size_t getSerializationSize() const = 0; + // Serialize the layer config to buffer. // TensorRT will call this func to serialize the configuration of TensorRT // engine. It should not be called by users. - virtual void serialize(void* buffer) = 0; + virtual void serialize(void* buffer) const = 0; - // Check format support. The default is FLOAT32 and NCHW. - bool supportsFormat(nvinfer1::DataType type, - nvinfer1::PluginFormat format) const override; - // Configure the layer - void configureWithFormat(const nvinfer1::Dims* input_dims, int num_inputs, - const nvinfer1::Dims* output_dims, int num_outputs, - nvinfer1::DataType type, - nvinfer1::PluginFormat format, - int max_batch_size) override; + void destroy() override { delete this; } + + virtual nvinfer1::IPluginV2Ext* clone() const = 0; + + void setPluginNamespace(const char* plugin_namespace) override { + namespace_ = plugin_namespace; + } + + const char* getPluginNamespace() const override { return namespace_.c_str(); } + + // IPluginV2Ext + nvinfer1::DataType getOutputDataType(int32_t index, + const nvinfer1::DataType* input_types, + int32_t nb_inputs) const override; + + bool isOutputBroadcastAcrossBatch(int32_t output_index, + const bool* input_is_broadcasted, + int32_t nb_inputs) const override; + + bool canBroadcastInputAcrossBatch(int32_t input_index) const override; + + void configurePlugin(const nvinfer1::Dims* input_dims, int32_t nb_inputs, + const nvinfer1::Dims* output_dims, int32_t nb_outputs, + const nvinfer1::DataType* input_types, + const nvinfer1::DataType* output_types, + const bool* input_is_broadcast, + const bool* output_is_broadcast, + nvinfer1::PluginFormat float_format, + int32_t max_batch_size) override; protected: // Deserialize input_dims, max_batch_size, data_type, data_format void deserializeBase(void const*& serial_data, // NOLINT size_t& serial_length); // NOLINT - size_t getBaseSerializationSize(); + size_t getBaseSerializationSize() const; // Serialize input_dims, max_batch_size, data_type, data_format - void serializeBase(void*& buffer); // NOLINT + void serializeBase(void*& buffer) const; // NOLINT std::vector input_dims_; size_t max_batch_size_; @@ -118,6 +156,9 @@ class PluginTensorRT : public nvinfer1::IPluginExt { std::vector inputs_; bool with_fp16_; + + private: + std::string namespace_; }; // TensorRT introduced IPluginV2Ext after 5.1, Paddle no longer supports From 5d63c78341e3e57e00debfcd0703448f779ab5b4 Mon Sep 17 00:00:00 2001 From: Reese Wang Date: Sun, 20 Jun 2021 16:56:56 +0800 Subject: [PATCH 03/46] trt plugin factory is not usable in IPluginV2 --- paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.cc | 2 ++ paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h | 2 ++ 2 files changed, 4 insertions(+) diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.cc b/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.cc index dd4e06ee2a900..94f5beccd56cb 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.cc +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.cc @@ -21,6 +21,7 @@ namespace inference { namespace tensorrt { namespace plugin { +#if false PluginTensorRT* PluginFactoryTensorRT::createPlugin(const char* layer_name, const void* serial_data, size_t serial_length) { @@ -45,6 +46,7 @@ bool PluginFactoryTensorRT::RegisterPlugin( } void PluginFactoryTensorRT::DestroyPlugins() { owned_plugins_.clear(); } +#endif } // namespace plugin } // namespace tensorrt diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h b/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h index 076dfbcf8f095..98ea916b870c9 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h @@ -33,6 +33,7 @@ namespace inference { namespace tensorrt { namespace plugin { +#if false class PluginFactoryTensorRT : public nvinfer1::IPluginFactory, public DeleteHelper { public: @@ -72,6 +73,7 @@ class TrtPluginRegistrar { trt_plugin_registrar##ctr UNUSED = \ paddle::inference::tensorrt::plugin::TrtPluginRegistrar( \ name, deserialize_func) +#endif } // namespace plugin } // namespace tensorrt From f7e94a0723b147ea3a53ef9fdff0552997d0848e Mon Sep 17 00:00:00 2001 From: Reese Wang Date: Sun, 20 Jun 2021 16:57:55 +0800 Subject: [PATCH 04/46] upgrade add plugin api to use IPluginV2 --- paddle/fluid/inference/tensorrt/engine.cc | 4 ++-- paddle/fluid/inference/tensorrt/engine.h | 7 +++---- 2 files changed, 5 insertions(+), 6 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/engine.cc b/paddle/fluid/inference/tensorrt/engine.cc index 99549fd6b5cbf..227d0abba3c1b 100644 --- a/paddle/fluid/inference/tensorrt/engine.cc +++ b/paddle/fluid/inference/tensorrt/engine.cc @@ -346,11 +346,11 @@ float *TensorRTEngine::GetWeightCPUData(const std::string &name, int TensorRTEngine::GetRuntimeBatch() { return runtime_batch_; } -nvinfer1::IPluginLayer *TensorRTEngine::AddPlugin( +nvinfer1::IPluginV2Layer *TensorRTEngine::AddPlugin( nvinfer1::ITensor *const *inputs, int num_inputs, plugin::PluginTensorRT *plugin) { owned_plugin_.emplace_back(plugin); - return network()->addPluginExt(inputs, num_inputs, *plugin); + return network()->addPluginV2(inputs, num_inputs, *plugin); } nvinfer1::IPluginV2Layer *TensorRTEngine::AddPluginV2Ext( diff --git a/paddle/fluid/inference/tensorrt/engine.h b/paddle/fluid/inference/tensorrt/engine.h index 7e5707269782e..6e31fda770d02 100644 --- a/paddle/fluid/inference/tensorrt/engine.h +++ b/paddle/fluid/inference/tensorrt/engine.h @@ -290,8 +290,7 @@ class TensorRTEngine { #endif } else { infer_engine_.reset(runtime->deserializeCudaEngine( - engine_serialized_data.c_str(), engine_serialized_data.size(), - &inference::Singleton::Global())); + engine_serialized_data.c_str(), engine_serialized_data.size())); } PADDLE_ENFORCE_NOT_NULL( infer_engine_, @@ -314,8 +313,8 @@ class TensorRTEngine { int GetDeviceId() { return device_id_; } - nvinfer1::IPluginLayer* AddPlugin(nvinfer1::ITensor* const* inputs, - int num_inputs, plugin::PluginTensorRT*); + nvinfer1::IPluginV2Layer* AddPlugin(nvinfer1::ITensor* const* inputs, + int num_inputs, plugin::PluginTensorRT*); nvinfer1::IPluginV2Layer* AddPluginV2Ext(nvinfer1::ITensor* const* inputs, int num_inputs, From ac6bd527f78e53057e7ed1e13f844488bd9872a0 Mon Sep 17 00:00:00 2001 From: Reese Wang Date: Sun, 20 Jun 2021 16:59:39 +0800 Subject: [PATCH 05/46] remove IPlugin register and adapt getSerializeSize(), serialize() --- .../tensorrt/plugin/elementwise_op_plugin.cu | 2 ++ .../tensorrt/plugin/elementwise_op_plugin.h | 6 +++--- .../inference/tensorrt/plugin/gelu_op_plugin.cu | 2 ++ .../inference/tensorrt/plugin/gelu_op_plugin.h | 5 ++--- .../tensorrt/plugin/hard_swish_op_plugin.cu | 2 ++ .../tensorrt/plugin/hard_swish_op_plugin.h | 14 +++++++------- .../tensorrt/plugin/instance_norm_op_plugin.cu | 2 ++ .../tensorrt/plugin/instance_norm_op_plugin.h | 7 +++---- .../tensorrt/plugin/layer_norm_op_plugin.cu | 2 ++ .../tensorrt/plugin/layer_norm_op_plugin.h | 7 +++---- .../inference/tensorrt/plugin/pool_op_plugin.cu | 2 ++ .../inference/tensorrt/plugin/pool_op_plugin.h | 7 +++---- .../inference/tensorrt/plugin/prelu_op_plugin.cu | 2 ++ .../inference/tensorrt/plugin/prelu_op_plugin.h | 7 +++---- .../inference/tensorrt/plugin/slice_op_plugin.cu | 6 ++++-- .../inference/tensorrt/plugin/slice_op_plugin.h | 5 ++--- .../inference/tensorrt/plugin/swish_op_plugin.cu | 2 ++ .../inference/tensorrt/plugin/swish_op_plugin.h | 7 +++---- 18 files changed, 49 insertions(+), 38 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu index cc17f8aa24817..c3f7e43260ffa 100644 --- a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu @@ -21,11 +21,13 @@ namespace inference { namespace tensorrt { namespace plugin { +#if false ElementWisePlugin *CreateElementWisePluginDeserialize(const void *buffer, size_t length) { return new ElementWisePlugin(buffer, length); } REGISTER_TRT_PLUGIN("elementwise_plugin", CreateElementWisePluginDeserialize); +#endif namespace details { template diff --git a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h index 75a1dd85f0f2c..5af89f218106b 100644 --- a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h @@ -62,14 +62,13 @@ class ElementWisePlugin : public PluginTensorRT { int enqueue(int batch_size, const void* const* inputs, void** outputs, void* workspace, cudaStream_t stream); - protected: - size_t getSerializationSize() override { + size_t getSerializationSize() const override { return SerializedSize(getPluginType()) + SerializedSize(axis_) + SerializedSize(dims_x_) + SerializedSize(dims_y_) + getBaseSerializationSize(); } - void serialize(void* buffer) override { + void serialize(void* buffer) const override { SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, type_.c_str()); @@ -78,6 +77,7 @@ class ElementWisePlugin : public PluginTensorRT { SerializeValue(&buffer, dims_y_); } + protected: std::string type_; nvinfer1::Dims dims_x_; nvinfer1::Dims dims_y_; diff --git a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu index deda2e2cc7247..852fbeb1f69c6 100644 --- a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu @@ -31,11 +31,13 @@ static const float kAT = 0.5; static const float kBT = 0.7978845608028654; // sqrt(2.0/M_PI) static const float kCT = 0.035677408136300125; // 0.044715 * sqrt(2.0/M_PI) +#if false GeluPlugin* CreateGeluPluginDeserialize(const void* buffer, size_t length) { return new GeluPlugin(buffer, length); } REGISTER_TRT_PLUGIN("gelu_plugin", CreateGeluPluginDeserialize); +#endif bool GeluPlugin::supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const { diff --git a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h index 23e507ee477e1..4ffaa0f8c7df1 100644 --- a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h @@ -47,14 +47,13 @@ class GeluPlugin : public PluginTensorRT { int enqueue(int batch_size, const void* const* inputs, void** outputs, void* workspace, cudaStream_t stream) override; - protected: - size_t getSerializationSize() override { + size_t getSerializationSize() const override { return getBaseSerializationSize() + SerializedSize(getPluginType()); } // TRT will call this func to serialize the configuration of TRT // It should not be called by users. - void serialize(void* buffer) override { + void serialize(void* buffer) const override { SerializeValue(&buffer, getPluginType()); serializeBase(buffer); } diff --git a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu index 8b2d0ac3cf70f..7edba358d0977 100644 --- a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu @@ -22,12 +22,14 @@ namespace inference { namespace tensorrt { namespace plugin { +#if false HardSwishPlugin* CreateHardSwishPluginDeserialize(const void* buffer, size_t length) { return new HardSwishPlugin(buffer, length); } REGISTER_TRT_PLUGIN("hard_swish_plugin", CreateHardSwishPluginDeserialize); +#endif nvinfer1::Dims HardSwishPlugin::getOutputDimensions( int index, const nvinfer1::Dims* in_dims, int nb_inputs) { diff --git a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h index 2e1e1d03baf7e..5c8c652be2e39 100644 --- a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h @@ -52,12 +52,7 @@ class HardSwishPlugin : public PluginTensorRT { int enqueue(int batchSize, const void* const* inputs, void** outputs, void* workspace, cudaStream_t stream) override; - protected: - float threshold_; - float scale_; - float offset_; - - size_t getSerializationSize() override { + size_t getSerializationSize() const override { return getBaseSerializationSize() + SerializedSize(threshold_) + SerializedSize(scale_) + SerializedSize(offset_) + SerializedSize(getPluginType()); @@ -65,13 +60,18 @@ class HardSwishPlugin : public PluginTensorRT { // TRT will call this func to serialize the configuration of TRT // It should not be called by users. - void serialize(void* buffer) override { + void serialize(void* buffer) const override { SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, threshold_); SerializeValue(&buffer, scale_); SerializeValue(&buffer, offset_); } + + protected: + float threshold_; + float scale_; + float offset_; }; } // namespace plugin diff --git a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu index a579743ee8ad1..f92cd18603fc6 100644 --- a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu @@ -40,12 +40,14 @@ cudnnStatus_t convert_trt2cudnn_dtype(nvinfer1::DataType trt_dtype, return CUDNN_STATUS_SUCCESS; } +#if false InstanceNormPlugin *CreateInstanceNormPluginDeserialize(const void *buffer, size_t length) { return new InstanceNormPlugin(buffer, length); } REGISTER_TRT_PLUGIN("instance_norm_plugin", CreateInstanceNormPluginDeserialize); +#endif int InstanceNormPlugin::initialize() { return 0; } diff --git a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h index 83422708f593d..d081dcb604923 100644 --- a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h @@ -38,8 +38,8 @@ class InstanceNormPlugin : public PluginTensorRT { cudnnHandle_t handle_; cudnnTensorDescriptor_t x_desc_, y_desc_, b_desc_; - protected: - size_t getSerializationSize() override { + public: + size_t getSerializationSize() const override { return getBaseSerializationSize() + SerializedSize(eps_) + SerializedSize(scale_) + SerializedSize(bias_) + SerializedSize(getPluginType()); @@ -48,7 +48,7 @@ class InstanceNormPlugin : public PluginTensorRT { // TRT will call this func when we need to serialize the configuration of // tensorrt. // It should not be called by users. - void serialize(void *buffer) override { + void serialize(void *buffer) const override { SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, eps_); @@ -56,7 +56,6 @@ class InstanceNormPlugin : public PluginTensorRT { SerializeValue(&buffer, bias_); } - public: explicit InstanceNormPlugin(const float eps, const std::vector scale, const std::vector bias) : eps_(eps), scale_(scale), bias_(bias) { diff --git a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu index f9341613a0f55..ef2a8bdee4275 100644 --- a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu @@ -25,11 +25,13 @@ namespace inference { namespace tensorrt { namespace plugin { +#if false LayerNormPlugin *CreateLayerNormPluginDeserialize(const void *buffer, size_t length) { return new LayerNormPlugin(buffer, length); } REGISTER_TRT_PLUGIN("layer_norm_plugin", CreateLayerNormPluginDeserialize); +#endif int LayerNormPlugin::initialize() { return 0; } diff --git a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h index 9c4c31b61e128..cb2ba114b5a5b 100644 --- a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h @@ -39,8 +39,8 @@ class LayerNormPlugin : public PluginTensorRT { std::vector mean_shape_; std::vector variance_shape_; - protected: - size_t getSerializationSize() override { + public: + size_t getSerializationSize() const override { return getBaseSerializationSize() + SerializedSize(bias_) + SerializedSize(scale_) + SerializedSize(begin_norm_axis_) + SerializedSize(eps_) + SerializedSize(mean_shape_) + @@ -50,7 +50,7 @@ class LayerNormPlugin : public PluginTensorRT { // TRT will call this func when we need to serialize the configuration of // tensorrt. // It should not be called by users. - void serialize(void* buffer) override { + void serialize(void* buffer) const override { SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, bias_); @@ -61,7 +61,6 @@ class LayerNormPlugin : public PluginTensorRT { SerializeValue(&buffer, variance_shape_); } - public: LayerNormPlugin(const float* bias, const int bias_num, const float* scale, const int scale_num, int begin_norm_axis, float eps, std::vector mean_shape, diff --git a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu index 154f61a2b7cd3..e57dcd63db4d3 100644 --- a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu @@ -21,10 +21,12 @@ namespace inference { namespace tensorrt { namespace plugin { +#if false PoolPlugin *CreatePoolPluginDeserialize(const void *buffer, size_t length) { return new PoolPlugin(buffer, length); } REGISTER_TRT_PLUGIN("pool_plugin", CreatePoolPluginDeserialize); +#endif nvinfer1::Dims PoolPlugin::getOutputDimensions(int index, const nvinfer1::Dims *inputDims, diff --git a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h index 6693a1fae4d43..504d4e97756dc 100644 --- a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h @@ -56,8 +56,8 @@ static std::vector CalcOutputSize(const std::vector& input_shape, } class PoolPlugin : public PluginTensorRT { - protected: - size_t getSerializationSize() override { + public: + size_t getSerializationSize() const override { return SerializedSize(getPluginType()) + SerializedSize(ceil_mode_) + SerializedSize(pool_type_) + SerializedSize(adaptive_) + SerializedSize(ksize_) + SerializedSize(strides_) + @@ -67,7 +67,7 @@ class PoolPlugin : public PluginTensorRT { // TRT will call this func when we need to serialize the configuration of // tensorrt. - void serialize(void* buffer) override { + void serialize(void* buffer) const override { SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, ceil_mode_); @@ -80,7 +80,6 @@ class PoolPlugin : public PluginTensorRT { SerializeValue(&buffer, output_shape_); } - public: enum class PoolType { max = 0, avg, diff --git a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu index 00182b87e984f..2f4f25ac283e4 100644 --- a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu @@ -27,10 +27,12 @@ namespace inference { namespace tensorrt { namespace plugin { +#if false PReluPlugin *CreatePreluPluginDeserialize(const void *buffer, size_t length) { return new PReluPlugin(buffer, length); } REGISTER_TRT_PLUGIN("prelu_plugin", CreatePreluPluginDeserialize); +#endif int PReluPlugin::initialize() { cudaMalloc(&p_gpu_weight_, sizeof(float) * weight_.size()); diff --git a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h index a0a24e70a01ef..013fc4d225f5c 100644 --- a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h @@ -33,8 +33,8 @@ class PReluPlugin : public PluginTensorRT { float* p_gpu_weight_; std::string mode_; - protected: - size_t getSerializationSize() override { + public: + size_t getSerializationSize() const override { return getBaseSerializationSize() + SerializedSize(mode_.c_str()) + SerializedSize(weight_) + SerializedSize(getPluginType()); } @@ -42,14 +42,13 @@ class PReluPlugin : public PluginTensorRT { // TRT will call this func when we need to serialize the configuration of // tensorrt. // It should not be called by users. - void serialize(void* buffer) override { + void serialize(void* buffer) const override { SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, weight_); SerializeValue(&buffer, mode_.c_str()); } - public: PReluPlugin(const float* weight, const int weight_num, std::string const& mode) : mode_(mode) { diff --git a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu index b44b3face92e1..50ee9d7ec1803 100644 --- a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu @@ -26,10 +26,12 @@ namespace inference { namespace tensorrt { namespace plugin { +#if false SlicePlugin *CreateSlicePluginDeserialize(const void *buffer, size_t length) { return new SlicePlugin(buffer, length); } REGISTER_TRT_PLUGIN("slice_plugin", CreateSlicePluginDeserialize); +#endif template __global__ void SliceKernel(int num, int dims, const T *input, @@ -188,13 +190,13 @@ int SlicePlugin::enqueue(int batch_size, const void *const *inputs, return cudaGetLastError() != cudaSuccess; } -size_t SlicePlugin::getSerializationSize() { +size_t SlicePlugin::getSerializationSize() const { return getBaseSerializationSize() + SerializedSize(getPluginType()) + SerializedSize(starts_) + SerializedSize(ends_) + SerializedSize(axes_); } -void SlicePlugin::serialize(void *buffer) { +void SlicePlugin::serialize(void *buffer) const { SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, starts_); diff --git a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h index 9d4f9a35c3b6f..7d15ea928a3a0 100644 --- a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h @@ -47,12 +47,11 @@ class SlicePlugin : public PluginTensorRT { int enqueue(int batch_size, const void* const* inputs, void** outputs, void* workspace, cudaStream_t stream) override; - protected: - size_t getSerializationSize() override; + size_t getSerializationSize() const override; // TRT will call this func to serialize the configuration of TRT // It should not be called by users. - void serialize(void* buffer) override; + void serialize(void* buffer) const override; private: std::vector starts_; diff --git a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu index 3847d999446e9..d13879443ff19 100644 --- a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu @@ -24,10 +24,12 @@ namespace inference { namespace tensorrt { namespace plugin { +#if false SwishPlugin *CreateSwishPluginDeserialize(const void *buffer, size_t length) { return new SwishPlugin(buffer, length); } REGISTER_TRT_PLUGIN("swish_plugin", CreateSwishPluginDeserialize); +#endif int SwishPlugin::initialize() { return 0; } diff --git a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h index 11579aadcc457..33f994d817d34 100644 --- a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h @@ -30,8 +30,8 @@ class SwishPlugin : public PluginTensorRT { private: float beta_; - protected: - size_t getSerializationSize() override { + public: + size_t getSerializationSize() const override { return SerializedSize(getPluginType()) + getBaseSerializationSize() + SerializedSize(beta_); } @@ -39,13 +39,12 @@ class SwishPlugin : public PluginTensorRT { // TRT will call this func when we need to serialize the configuration of // tensorrt. // It should not be called by users. - void serialize(void* buffer) override { + void serialize(void* buffer) const override { SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, beta_); } - public: explicit SwishPlugin(const float beta, const bool with_fp16) : beta_(beta) { with_fp16_ = with_fp16; } From b311c3f956c7f28edcac4099683e1d8a9cce97b6 Mon Sep 17 00:00:00 2001 From: Reese Wang Date: Sun, 20 Jun 2021 17:00:15 +0800 Subject: [PATCH 06/46] adapt IPluginV2Layer --- paddle/fluid/inference/tensorrt/convert/elementwise_op.cc | 2 +- paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc b/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc index df2400854414c..6820cc8cda668 100644 --- a/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc +++ b/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc @@ -253,7 +253,7 @@ class ElementwiseTensorOpConverter : public OpConverter { new plugin::ElementWisePlugin(op_type_, dims_x, dims_y, axis); plugin->AddInput(X); plugin->AddInput(Y); - nvinfer1::IPluginLayer* plugin_layer = engine_->AddPlugin( + auto* plugin_layer = engine_->AddPlugin( plugin->GetInputs().data(), 2, reinterpret_cast(plugin)); diff --git a/paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc b/paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc index 2fd0d82bb1ea3..b7097fc05680d 100644 --- a/paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc +++ b/paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc @@ -74,7 +74,7 @@ class InstanceNormOpConverter : public OpConverter { plugin::InstanceNormPlugin* plugin = new plugin::InstanceNormPlugin(eps, scale_v, bias_v); plugin->getPluginType(); - nvinfer1::IPluginLayer* layer = engine_->AddPlugin(&input, 1, plugin); + auto* layer = engine_->AddPlugin(&input, 1, plugin); auto output_name = op_desc.Output("Y")[0]; RreplenishLayerAndOutput(layer, "instance_norm", {output_name}, test_mode); From 99e146107b68badf62fbd2458741ebf26ccb3406 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Thu, 24 Jun 2021 07:48:13 +0000 Subject: [PATCH 07/46] downgrade to IPluginV2 --- .../tensorrt/plugin/elementwise_op_plugin.cu | 8 ---- .../tensorrt/plugin/gelu_op_plugin.cu | 8 ---- .../tensorrt/plugin/hard_swish_op_plugin.cu | 9 ---- .../plugin/instance_norm_op_plugin.cu | 9 ---- .../tensorrt/plugin/layer_norm_op_plugin.cu | 8 ---- .../tensorrt/plugin/pool_op_plugin.cu | 7 --- .../tensorrt/plugin/prelu_op_plugin.cu | 7 --- .../tensorrt/plugin/slice_op_plugin.cu | 7 --- .../tensorrt/plugin/swish_op_plugin.cu | 7 --- .../inference/tensorrt/plugin/trt_plugin.cc | 4 +- .../inference/tensorrt/plugin/trt_plugin.h | 46 +++++++++---------- 11 files changed, 25 insertions(+), 95 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu index c3f7e43260ffa..1a2e111563cae 100644 --- a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu @@ -21,14 +21,6 @@ namespace inference { namespace tensorrt { namespace plugin { -#if false -ElementWisePlugin *CreateElementWisePluginDeserialize(const void *buffer, - size_t length) { - return new ElementWisePlugin(buffer, length); -} -REGISTER_TRT_PLUGIN("elementwise_plugin", CreateElementWisePluginDeserialize); -#endif - namespace details { template struct Add { diff --git a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu index 852fbeb1f69c6..8bf4c5af700b2 100644 --- a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu @@ -31,14 +31,6 @@ static const float kAT = 0.5; static const float kBT = 0.7978845608028654; // sqrt(2.0/M_PI) static const float kCT = 0.035677408136300125; // 0.044715 * sqrt(2.0/M_PI) -#if false -GeluPlugin* CreateGeluPluginDeserialize(const void* buffer, size_t length) { - return new GeluPlugin(buffer, length); -} - -REGISTER_TRT_PLUGIN("gelu_plugin", CreateGeluPluginDeserialize); -#endif - bool GeluPlugin::supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const { if (with_fp16_) { diff --git a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu index 7edba358d0977..eb3504d500ab2 100644 --- a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu @@ -22,15 +22,6 @@ namespace inference { namespace tensorrt { namespace plugin { -#if false -HardSwishPlugin* CreateHardSwishPluginDeserialize(const void* buffer, - size_t length) { - return new HardSwishPlugin(buffer, length); -} - -REGISTER_TRT_PLUGIN("hard_swish_plugin", CreateHardSwishPluginDeserialize); -#endif - nvinfer1::Dims HardSwishPlugin::getOutputDimensions( int index, const nvinfer1::Dims* in_dims, int nb_inputs) { assert(nb_inputs == 1); diff --git a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu index f92cd18603fc6..41ca867e01d44 100644 --- a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu @@ -40,15 +40,6 @@ cudnnStatus_t convert_trt2cudnn_dtype(nvinfer1::DataType trt_dtype, return CUDNN_STATUS_SUCCESS; } -#if false -InstanceNormPlugin *CreateInstanceNormPluginDeserialize(const void *buffer, - size_t length) { - return new InstanceNormPlugin(buffer, length); -} -REGISTER_TRT_PLUGIN("instance_norm_plugin", - CreateInstanceNormPluginDeserialize); -#endif - int InstanceNormPlugin::initialize() { return 0; } nvinfer1::Dims InstanceNormPlugin::getOutputDimensions( diff --git a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu index ef2a8bdee4275..fce9bcb2f675a 100644 --- a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu @@ -25,14 +25,6 @@ namespace inference { namespace tensorrt { namespace plugin { -#if false -LayerNormPlugin *CreateLayerNormPluginDeserialize(const void *buffer, - size_t length) { - return new LayerNormPlugin(buffer, length); -} -REGISTER_TRT_PLUGIN("layer_norm_plugin", CreateLayerNormPluginDeserialize); -#endif - int LayerNormPlugin::initialize() { return 0; } nvinfer1::Dims LayerNormPlugin::getOutputDimensions( diff --git a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu index e57dcd63db4d3..5156d0f8ee39d 100644 --- a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu @@ -21,13 +21,6 @@ namespace inference { namespace tensorrt { namespace plugin { -#if false -PoolPlugin *CreatePoolPluginDeserialize(const void *buffer, size_t length) { - return new PoolPlugin(buffer, length); -} -REGISTER_TRT_PLUGIN("pool_plugin", CreatePoolPluginDeserialize); -#endif - nvinfer1::Dims PoolPlugin::getOutputDimensions(int index, const nvinfer1::Dims *inputDims, int nbInputs) { diff --git a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu index 2f4f25ac283e4..54db6829b6c67 100644 --- a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu @@ -27,13 +27,6 @@ namespace inference { namespace tensorrt { namespace plugin { -#if false -PReluPlugin *CreatePreluPluginDeserialize(const void *buffer, size_t length) { - return new PReluPlugin(buffer, length); -} -REGISTER_TRT_PLUGIN("prelu_plugin", CreatePreluPluginDeserialize); -#endif - int PReluPlugin::initialize() { cudaMalloc(&p_gpu_weight_, sizeof(float) * weight_.size()); cudaMemcpy(p_gpu_weight_, weight_.data(), weight_.size() * sizeof(float), diff --git a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu index 50ee9d7ec1803..885d123975302 100644 --- a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu @@ -26,13 +26,6 @@ namespace inference { namespace tensorrt { namespace plugin { -#if false -SlicePlugin *CreateSlicePluginDeserialize(const void *buffer, size_t length) { - return new SlicePlugin(buffer, length); -} -REGISTER_TRT_PLUGIN("slice_plugin", CreateSlicePluginDeserialize); -#endif - template __global__ void SliceKernel(int num, int dims, const T *input, const int *offsets_info, T *output) { diff --git a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu index d13879443ff19..66095bcb91117 100644 --- a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu @@ -24,13 +24,6 @@ namespace inference { namespace tensorrt { namespace plugin { -#if false -SwishPlugin *CreateSwishPluginDeserialize(const void *buffer, size_t length) { - return new SwishPlugin(buffer, length); -} -REGISTER_TRT_PLUGIN("swish_plugin", CreateSwishPluginDeserialize); -#endif - int SwishPlugin::initialize() { return 0; } nvinfer1::Dims SwishPlugin::getOutputDimensions(int index, diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc index 751deac45859b..8e9e74c9c6715 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc @@ -80,7 +80,7 @@ void PluginTensorRT::configureWithFormat( input_dims_.assign(input_dims, input_dims + num_inputs); max_batch_size_ = max_batch_size; } - +/* nvinfer1::DataType PluginTensorRT::getOutputDataType( int32_t index, const nvinfer1::DataType* input_types, int32_t nb_inputs) const { @@ -104,7 +104,7 @@ void PluginTensorRT::configurePlugin( const nvinfer1::DataType* output_types, const bool* input_is_broadcast, const bool* output_is_broadcast, nvinfer1::PluginFormat float_format, int32_t max_batch_size) {} - +*/ void PluginTensorRTV2Ext::serializeBase(void*& buffer) const { Seria(buffer, input_dims_, max_batch_size_, data_type_, data_format_, with_fp16_); diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h index 428e1dfc50eb8..f9c57952150fb 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h @@ -45,7 +45,7 @@ typedef std::function typedef std::function PluginConstructFunc; // Deprecated. Do not inherit this class, please refer to PluginTensorRTV2Ext -class PluginTensorRT : public nvinfer1::IPluginV2Ext { +class PluginTensorRT : public nvinfer1::IPluginV2 { public: PluginTensorRT() : with_fp16_(false) {} // It was used for TensorRT deserialization. @@ -113,34 +113,34 @@ class PluginTensorRT : public nvinfer1::IPluginV2Ext { void destroy() override { delete this; } - virtual nvinfer1::IPluginV2Ext* clone() const = 0; + virtual nvinfer1::IPluginV2* clone() const = 0; void setPluginNamespace(const char* plugin_namespace) override { namespace_ = plugin_namespace; } const char* getPluginNamespace() const override { return namespace_.c_str(); } - - // IPluginV2Ext - nvinfer1::DataType getOutputDataType(int32_t index, - const nvinfer1::DataType* input_types, - int32_t nb_inputs) const override; - - bool isOutputBroadcastAcrossBatch(int32_t output_index, - const bool* input_is_broadcasted, - int32_t nb_inputs) const override; - - bool canBroadcastInputAcrossBatch(int32_t input_index) const override; - - void configurePlugin(const nvinfer1::Dims* input_dims, int32_t nb_inputs, - const nvinfer1::Dims* output_dims, int32_t nb_outputs, - const nvinfer1::DataType* input_types, - const nvinfer1::DataType* output_types, - const bool* input_is_broadcast, - const bool* output_is_broadcast, - nvinfer1::PluginFormat float_format, - int32_t max_batch_size) override; - + /* + // IPluginV2Ext + nvinfer1::DataType getOutputDataType(int32_t index, + const nvinfer1::DataType* input_types, + int32_t nb_inputs) const override; + + bool isOutputBroadcastAcrossBatch(int32_t output_index, + const bool* input_is_broadcasted, + int32_t nb_inputs) const override; + + bool canBroadcastInputAcrossBatch(int32_t input_index) const override; + + void configurePlugin(const nvinfer1::Dims* input_dims, int32_t nb_inputs, + const nvinfer1::Dims* output_dims, int32_t nb_outputs, + const nvinfer1::DataType* input_types, + const nvinfer1::DataType* output_types, + const bool* input_is_broadcast, + const bool* output_is_broadcast, + nvinfer1::PluginFormat float_format, + int32_t max_batch_size) override; + */ protected: // Deserialize input_dims, max_batch_size, data_type, data_format void deserializeBase(void const*& serial_data, // NOLINT From 686b368ec2a35afd0935736119286ae01485679d Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Thu, 24 Jun 2021 10:17:39 +0000 Subject: [PATCH 08/46] implement elementwise clone --- paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h index ab25ad8a36cdb..bd0be4d171778 100644 --- a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h @@ -46,8 +46,7 @@ class ElementWisePlugin : public PluginTensorRT { } ElementWisePlugin* clone() const override { - // return new ElementWisePlugin(dims_x_, dims_y_, axis_); - return nullptr; + return new ElementWisePlugin(type_, dims_x_, dims_y_, axis_); } const char* getPluginType() const override { return "elementwise_plugin"; } From fffb0b509e7a287c0da6c248f792e67be072d881 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 03:59:15 +0000 Subject: [PATCH 09/46] add gelu plugin creator and fix gelu serialization bug --- .../tensorrt/plugin/gelu_op_plugin.h | 45 +++++++++++++++++-- 1 file changed, 41 insertions(+), 4 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h index 6a4c2ce702016..d03306bb9bc39 100644 --- a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h @@ -52,17 +52,54 @@ class GeluPlugin : public PluginTensorRT { void* workspace, cudaStream_t stream) override; size_t getSerializationSize() const override { - return getBaseSerializationSize() + SerializedSize(getPluginType()); + return getBaseSerializationSize(); } // TRT will call this func to serialize the configuration of TRT // It should not be called by users. void serialize(void* buffer) const override { - SerializeValue(&buffer, getPluginType()); serializeBase(buffer); } }; +class GeluPluginCreator : public nvinfer1::IPluginCreator { + public: + GeluPluginCreator() {} + const char* getPluginName() const override { return "gelu_plugin"; } + + const char* getPluginVersion() const override { return "1"; } + + const nvinfer1::PluginFieldCollection* getFieldNames() override { + return &field_collection_; + } + + nvinfer1::IPluginV2* createPlugin( + const char* name, const nvinfer1::PluginFieldCollection* fc) override { + return nullptr; + } + + nvinfer1::IPluginV2* deserializePlugin(const char* name, + const void* serial_data, + size_t serial_length) override { + return new GeluPlugin(serial_data, serial_length); + } + + void setPluginNamespace(const char* lib_namespace) override { + plugin_namespace_ = lib_namespace; + } + + const char* getPluginNamespace() const override { + return plugin_namespace_.c_str(); + } + + private: + std::string plugin_namespace_; + std::string plugin_name_; + nvinfer1::PluginFieldCollection field_collection_{0, nullptr}; + std::vector plugin_attributes_; +}; +REGISTER_TRT_PLUGIN_V2(GeluPluginCreator); + #if IS_TRT_VERSION_GE(6000) class GeluPluginDynamic : public DynamicPluginTensorRT { public: @@ -76,7 +113,7 @@ class GeluPluginDynamic : public DynamicPluginTensorRT { return new GeluPluginDynamic(with_fp16_); } - const char* getPluginType() const override { return "gelu_plugin"; } + const char* getPluginType() const override { return "gelu_plugin_dynamic"; } int getNbOutputs() const override { return 1; } int initialize() override { return 0; } @@ -121,7 +158,7 @@ class GeluPluginDynamic : public DynamicPluginTensorRT { class GeluPluginDynamicCreator : public nvinfer1::IPluginCreator { public: GeluPluginDynamicCreator() {} - const char* getPluginName() const override { return "gelu_plugin"; } + const char* getPluginName() const override { return "gelu_plugin_dynamic"; } const char* getPluginVersion() const override { return "1"; } From b55ac3aa05db33914d6e800495588988a2f549de Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 03:59:51 +0000 Subject: [PATCH 10/46] add swish plugin creator and fix swish serialization bug --- .../tensorrt/plugin/swish_op_plugin.h | 46 +++++++++++++++++-- 1 file changed, 42 insertions(+), 4 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h index d30b16477c8f1..4a181718b8578 100644 --- a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h @@ -32,7 +32,7 @@ class SwishPlugin : public PluginTensorRT { public: size_t getSerializationSize() const override { - return SerializedSize(getPluginType()) + getBaseSerializationSize() + + return getBaseSerializationSize() + SerializedSize(beta_); } @@ -40,7 +40,6 @@ class SwishPlugin : public PluginTensorRT { // tensorrt. // It should not be called by users. void serialize(void* buffer) const override { - SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, beta_); } @@ -74,6 +73,45 @@ class SwishPlugin : public PluginTensorRT { void* workspace, cudaStream_t stream) override; }; +class SwishPluginCreator : public nvinfer1::IPluginCreator { + public: + SwishPluginCreator() {} + const char* getPluginName() const override { return "swish_plugin"; } + + const char* getPluginVersion() const override { return "1"; } + + const nvinfer1::PluginFieldCollection* getFieldNames() override { + return &field_collection_; + } + + nvinfer1::IPluginV2* createPlugin( + const char* name, const nvinfer1::PluginFieldCollection* fc) override { + return nullptr; + } + + nvinfer1::IPluginV2* deserializePlugin(const char* name, + const void* serial_data, + size_t serial_length) override { + auto plugin = new SwishPlugin(serial_data, serial_length); + return plugin; + } + + void setPluginNamespace(const char* lib_namespace) override { + plugin_namespace_ = lib_namespace; + } + + const char* getPluginNamespace() const override { + return plugin_namespace_.c_str(); + } + + private: + std::string plugin_namespace_; + std::string plugin_name_; + nvinfer1::PluginFieldCollection field_collection_{0, nullptr}; + std::vector plugin_attributes_; +}; +REGISTER_TRT_PLUGIN_V2(SwishPluginCreator); + #if IS_TRT_VERSION_GE(6000) class SwishPluginDynamic : public DynamicPluginTensorRT { public: @@ -89,7 +127,7 @@ class SwishPluginDynamic : public DynamicPluginTensorRT { return new SwishPluginDynamic(beta_, with_fp16_); } - const char* getPluginType() const override { return "swish_plugin"; } + const char* getPluginType() const override { return "swish_plugin_dynamic"; } int getNbOutputs() const override { return 1; } int initialize() override; @@ -133,7 +171,7 @@ class SwishPluginDynamic : public DynamicPluginTensorRT { class SwishPluginDynamicCreator : public nvinfer1::IPluginCreator { public: SwishPluginDynamicCreator() {} - const char* getPluginName() const override { return "swish_plugin"; } + const char* getPluginName() const override { return "swish_plugin_dynamic"; } const char* getPluginVersion() const override { return "1"; } From c5b08bb7e842e65df98fdd0af860a36267e142c3 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 05:03:00 +0000 Subject: [PATCH 11/46] format --- paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h | 4 +--- paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h | 3 +-- 2 files changed, 2 insertions(+), 5 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h index d03306bb9bc39..1366a299e8f41 100644 --- a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h @@ -57,9 +57,7 @@ class GeluPlugin : public PluginTensorRT { // TRT will call this func to serialize the configuration of TRT // It should not be called by users. - void serialize(void* buffer) const override { - serializeBase(buffer); - } + void serialize(void* buffer) const override { serializeBase(buffer); } }; class GeluPluginCreator : public nvinfer1::IPluginCreator { diff --git a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h index 4a181718b8578..1b8342c8c895c 100644 --- a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h @@ -32,8 +32,7 @@ class SwishPlugin : public PluginTensorRT { public: size_t getSerializationSize() const override { - return getBaseSerializationSize() + - SerializedSize(beta_); + return getBaseSerializationSize() + SerializedSize(beta_); } // TRT will call this func when we need to serialize the configuration of From 326e093d050108fa984ba7dd86b0cb54f316ad36 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 10:34:00 +0000 Subject: [PATCH 12/46] fix typo --- paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc b/paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc index 1da44c98f36a0..976fe9502acd6 100644 --- a/paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc +++ b/paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc @@ -61,7 +61,8 @@ class ShuffleChannelOpConverter : public OpConverter { reshape_layer->setReshapeDimensions(reshape_dim2); auto output_name = op_desc.Output("Out")[0]; - RreplenishLayerAndOutput(reshape_layer, "concat", {output_name}, test_mode); + RreplenishLayerAndOutput(reshape_layer, "shuffle_channel", {output_name}, + test_mode); } }; From d43425235a6943e62197b5ac9f4abd111be59747 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 10:34:51 +0000 Subject: [PATCH 13/46] add elementwise plugin creator and fix serialization --- .../tensorrt/plugin/elementwise_op_plugin.h | 38 +++++++++++++++---- 1 file changed, 31 insertions(+), 7 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h index bd0be4d171778..5dd3142c75839 100644 --- a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h @@ -40,9 +40,12 @@ class ElementWisePlugin : public PluginTensorRT { const char* elementwise_type; DeserializeValue(&serial_data, &serial_length, &elementwise_type); type_ = std::string(elementwise_type); - DeserializeValue(&serial_data, &serial_length, &axis_); DeserializeValue(&serial_data, &serial_length, &dims_x_); DeserializeValue(&serial_data, &serial_length, &dims_y_); + DeserializeValue(&serial_data, &serial_length, &axis_); + DeserializeValue(&serial_data, &serial_length, &prev_size_); + DeserializeValue(&serial_data, &serial_length, &midd_size_); + DeserializeValue(&serial_data, &serial_length, &post_size_); } ElementWisePlugin* clone() const override { @@ -65,18 +68,21 @@ class ElementWisePlugin : public PluginTensorRT { void* workspace, cudaStream_t stream); size_t getSerializationSize() const override { - return SerializedSize(getPluginType()) + SerializedSize(axis_) + + return getBaseSerializationSize() + SerializedSize(type_.c_str()) + SerializedSize(dims_x_) + SerializedSize(dims_y_) + - getBaseSerializationSize(); + SerializedSize(axis_) + SerializedSize(prev_size_) + + SerializedSize(midd_size_) + SerializedSize(post_size_); } void serialize(void* buffer) const override { - SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, type_.c_str()); - SerializeValue(&buffer, axis_); SerializeValue(&buffer, dims_x_); SerializeValue(&buffer, dims_y_); + SerializeValue(&buffer, axis_); + SerializeValue(&buffer, prev_size_); + SerializeValue(&buffer, midd_size_); + SerializeValue(&buffer, post_size_); } protected: @@ -89,6 +95,20 @@ class ElementWisePlugin : public PluginTensorRT { int post_size_; }; +class ElementWisePluginCreator : public TensorRTPluginCreator { + public: + const char* getPluginName() const override { return "elementwise_plugin"; } + + const char* getPluginVersion() const override { return "1"; } + + nvinfer1::IPluginV2* deserializePlugin(const char* name, + const void* serial_data, + size_t serial_length) override { + return new ElementWisePlugin(serial_data, serial_length); + } +}; +REGISTER_TRT_PLUGIN_V2(ElementWisePluginCreator); + #if IS_TRT_VERSION_GE(6000) class ElementwisePluginDynamic : public DynamicPluginTensorRT { public: @@ -104,7 +124,9 @@ class ElementwisePluginDynamic : public DynamicPluginTensorRT { return new ElementwisePluginDynamic(type_, axis_); } - const char* getPluginType() const override { return "elementwise_plugin"; } + const char* getPluginType() const override { + return "elementwise_plugin_dynamic"; + } int getNbOutputs() const override { return 1; } int initialize() override; @@ -149,7 +171,9 @@ class ElementwisePluginDynamic : public DynamicPluginTensorRT { class ElementwisePluginDynamicCreator : public nvinfer1::IPluginCreator { public: ElementwisePluginDynamicCreator() {} - const char* getPluginName() const override { return "elementwise_plugin"; } + const char* getPluginName() const override { + return "elementwise_plugin_dynamic"; + } const char* getPluginVersion() const override { return "1"; } From 3d36eceaaff255d94b10d22fbacf2c9606abedfb Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 10:36:02 +0000 Subject: [PATCH 14/46] add base creator class --- .../inference/tensorrt/plugin/trt_plugin.cc | 17 +++++++++++ .../inference/tensorrt/plugin/trt_plugin.h | 28 +++++++++++++++++++ 2 files changed, 45 insertions(+) diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc index 8e9e74c9c6715..64cf94fc93357 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc @@ -134,6 +134,23 @@ void PluginTensorRTV2Ext::configurePlugin( data_type_ = input_types[0]; } +const nvinfer1::PluginFieldCollection* TensorRTPluginCreator::getFieldNames() { + return &field_collection_; +} + +nvinfer1::IPluginV2* TensorRTPluginCreator::createPlugin( + const char* name, const nvinfer1::PluginFieldCollection* fc) { + return nullptr; +} + +void TensorRTPluginCreator::setPluginNamespace(const char* lib_namespace) { + plugin_namespace_ = lib_namespace; +} + +const char* TensorRTPluginCreator::getPluginNamespace() const { + return plugin_namespace_.c_str(); +} + } // namespace plugin } // namespace tensorrt } // namespace inference diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h index 94027f3fab065..9cc99a4da3fa7 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h @@ -346,6 +346,34 @@ class DynamicPluginTensorRT : public nvinfer1::IPluginV2DynamicExt { }; #endif +class TensorRTPluginCreator : public nvinfer1::IPluginCreator { + public: + TensorRTPluginCreator() = default; + + virtual const char* getPluginName() const = 0; + + virtual const char* getPluginVersion() const = 0; + + const nvinfer1::PluginFieldCollection* getFieldNames() override; + + nvinfer1::IPluginV2* createPlugin( + const char* name, const nvinfer1::PluginFieldCollection* fc) override; + + virtual nvinfer1::IPluginV2* deserializePlugin(const char* name, + const void* serial_data, + size_t serial_length) = 0; + + void setPluginNamespace(const char* lib_namespace) override; + + const char* getPluginNamespace() const override; + + private: + std::string plugin_namespace_; + std::string plugin_name_; + nvinfer1::PluginFieldCollection field_collection_{0, nullptr}; + std::vector plugin_attributes_; +}; + template class TrtPluginRegistrarV2 { public: From 2020e58c5a66d31d07b20863477d97510f2f41b5 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 10:36:19 +0000 Subject: [PATCH 15/46] add gelu plugin creator --- .../tensorrt/plugin/gelu_op_plugin.h | 53 +------------------ 1 file changed, 2 insertions(+), 51 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h index 1366a299e8f41..6fdd9791a61bd 100644 --- a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h @@ -60,41 +60,17 @@ class GeluPlugin : public PluginTensorRT { void serialize(void* buffer) const override { serializeBase(buffer); } }; -class GeluPluginCreator : public nvinfer1::IPluginCreator { +class GeluPluginCreator : public TensorRTPluginCreator { public: - GeluPluginCreator() {} const char* getPluginName() const override { return "gelu_plugin"; } const char* getPluginVersion() const override { return "1"; } - const nvinfer1::PluginFieldCollection* getFieldNames() override { - return &field_collection_; - } - - nvinfer1::IPluginV2* createPlugin( - const char* name, const nvinfer1::PluginFieldCollection* fc) override { - return nullptr; - } - nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serial_data, size_t serial_length) override { return new GeluPlugin(serial_data, serial_length); } - - void setPluginNamespace(const char* lib_namespace) override { - plugin_namespace_ = lib_namespace; - } - - const char* getPluginNamespace() const override { - return plugin_namespace_.c_str(); - } - - private: - std::string plugin_namespace_; - std::string plugin_name_; - nvinfer1::PluginFieldCollection field_collection_{0, nullptr}; - std::vector plugin_attributes_; }; REGISTER_TRT_PLUGIN_V2(GeluPluginCreator); @@ -153,44 +129,19 @@ class GeluPluginDynamic : public DynamicPluginTensorRT { void destroy() override { delete this; } }; -class GeluPluginDynamicCreator : public nvinfer1::IPluginCreator { +class GeluPluginDynamicCreator : public TensorRTPluginCreator { public: - GeluPluginDynamicCreator() {} const char* getPluginName() const override { return "gelu_plugin_dynamic"; } const char* getPluginVersion() const override { return "1"; } - const nvinfer1::PluginFieldCollection* getFieldNames() override { - return &field_collection_; - } - - nvinfer1::IPluginV2* createPlugin( - const char* name, const nvinfer1::PluginFieldCollection* fc) override { - return nullptr; - } - nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serial_data, size_t serial_length) override { auto plugin = new GeluPluginDynamic(serial_data, serial_length); return plugin; } - - void setPluginNamespace(const char* lib_namespace) override { - plugin_namespace_ = lib_namespace; - } - - const char* getPluginNamespace() const override { - return plugin_namespace_.c_str(); - } - - private: - std::string plugin_namespace_; - std::string plugin_name_; - nvinfer1::PluginFieldCollection field_collection_{0, nullptr}; - std::vector plugin_attributes_; }; - REGISTER_TRT_PLUGIN_V2(GeluPluginDynamicCreator); #endif From 84e367564b2a55481138fb00da24e3740de43c40 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 10:36:59 +0000 Subject: [PATCH 16/46] add hard swish creator and fix serialization --- .../tensorrt/plugin/hard_swish_op_plugin.h | 18 +++++++++++++++--- 1 file changed, 15 insertions(+), 3 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h index 095d00e572edb..42c47959988a5 100644 --- a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h @@ -58,14 +58,12 @@ class HardSwishPlugin : public PluginTensorRT { size_t getSerializationSize() const override { return getBaseSerializationSize() + SerializedSize(threshold_) + - SerializedSize(scale_) + SerializedSize(offset_) + - SerializedSize(getPluginType()); + SerializedSize(scale_) + SerializedSize(offset_); } // TRT will call this func to serialize the configuration of TRT // It should not be called by users. void serialize(void* buffer) const override { - SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, threshold_); SerializeValue(&buffer, scale_); @@ -78,6 +76,20 @@ class HardSwishPlugin : public PluginTensorRT { float offset_; }; +class HardSwishPluginCreator : public TensorRTPluginCreator { + public: + const char* getPluginName() const override { return "hard_swish_plugin"; } + + const char* getPluginVersion() const override { return "1"; } + + nvinfer1::IPluginV2* deserializePlugin(const char* name, + const void* serial_data, + size_t serial_length) override { + return new HardSwishPlugin(serial_data, serial_length); + } +}; +REGISTER_TRT_PLUGIN_V2(HardSwishPluginCreator); + } // namespace plugin } // namespace tensorrt } // namespace inference From 20c91fa9d6f6a61293b6d8375a6c963d3ac2e6b4 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 11:26:57 +0000 Subject: [PATCH 17/46] add instance norm creator and fix serialization --- .../tensorrt/plugin/instance_norm_op_plugin.h | 18 +++++++++++++++--- 1 file changed, 15 insertions(+), 3 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h index c47209feef913..16f03c7e573c7 100644 --- a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h @@ -41,15 +41,13 @@ class InstanceNormPlugin : public PluginTensorRT { public: size_t getSerializationSize() const override { return getBaseSerializationSize() + SerializedSize(eps_) + - SerializedSize(scale_) + SerializedSize(bias_) + - SerializedSize(getPluginType()); + SerializedSize(scale_) + SerializedSize(bias_); } // TRT will call this func when we need to serialize the configuration of // tensorrt. // It should not be called by users. void serialize(void *buffer) const override { - SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, eps_); SerializeValue(&buffer, scale_); @@ -115,6 +113,20 @@ class InstanceNormPlugin : public PluginTensorRT { } }; +class InstanceNormPluginCreator : public TensorRTPluginCreator { + public: + const char *getPluginName() const override { return "instance_norm_plugin"; } + + const char *getPluginVersion() const override { return "1"; } + + nvinfer1::IPluginV2 *deserializePlugin(const char *name, + const void *serial_data, + size_t serial_length) override { + return new InstanceNormPlugin(serial_data, serial_length); + } +}; +REGISTER_TRT_PLUGIN_V2(InstanceNormPluginCreator); + } // namespace plugin } // namespace tensorrt } // namespace inference From 8fcd8fd5c7a98fe5c152ef76f036dbaf437f93f8 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 11:54:31 +0000 Subject: [PATCH 18/46] add layer norm creator and fix serialization --- .../tensorrt/plugin/layer_norm_op_plugin.h | 56 ++++++++----------- 1 file changed, 24 insertions(+), 32 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h index 45df62db71bb6..caa3c21db63fa 100644 --- a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h @@ -44,14 +44,13 @@ class LayerNormPlugin : public PluginTensorRT { return getBaseSerializationSize() + SerializedSize(bias_) + SerializedSize(scale_) + SerializedSize(begin_norm_axis_) + SerializedSize(eps_) + SerializedSize(mean_shape_) + - SerializedSize(variance_shape_) + SerializedSize(getPluginType()); + SerializedSize(variance_shape_); } // TRT will call this func when we need to serialize the configuration of // tensorrt. // It should not be called by users. void serialize(void* buffer) const override { - SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, bias_); SerializeValue(&buffer, scale_); @@ -95,7 +94,7 @@ class LayerNormPlugin : public PluginTensorRT { mean_shape_, variance_shape_); } - const char* getPluginType() const override { return "layer_norm_plugin"; } + const char* getPluginType() const override { return "layernorm_plugin"; } int getNbOutputs() const override { return 1; } nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) override; @@ -107,6 +106,20 @@ class LayerNormPlugin : public PluginTensorRT { void* workspace, cudaStream_t stream) override; }; +class LayerNormPluginCreator : public TensorRTPluginCreator { + public: + const char* getPluginName() const override { return "layernorm_plugin"; } + + const char* getPluginVersion() const override { return "1"; } + + nvinfer1::IPluginV2* deserializePlugin(const char* name, + const void* serial_data, + size_t serial_length) override { + return new LayerNormPlugin(serial_data, serial_length); + } +}; +REGISTER_TRT_PLUGIN_V2(LayerNormPluginCreator); + class LayerNormPluginDynamic : public DynamicPluginTensorRT { public: LayerNormPluginDynamic(const float* bias, const int bias_num, @@ -138,7 +151,9 @@ class LayerNormPluginDynamic : public DynamicPluginTensorRT { mean_shape_, variance_shape_); } - const char* getPluginType() const override { return "layernorm_plugin"; } + const char* getPluginType() const override { + return "layernorm_plugin_dynamic"; + } int getNbOutputs() const override { return 1; } int initialize() override { return 0; } @@ -200,42 +215,19 @@ class LayerNormPluginDynamic : public DynamicPluginTensorRT { std::vector variance_shape_; }; -class LayerNormPluginDynamicCreator : public nvinfer1::IPluginCreator { +class LayerNormPluginDynamicCreator : public TensorRTPluginCreator { public: - LayerNormPluginDynamicCreator() {} - const char* getPluginName() const override { return "layernorm_plugin"; } - - const char* getPluginVersion() const override { return "1"; } - - const nvinfer1::PluginFieldCollection* getFieldNames() override { - return &field_collection_; + const char* getPluginName() const override { + return "layernorm_plugin_dynamic"; } - nvinfer1::IPluginV2* createPlugin( - const char* name, const nvinfer1::PluginFieldCollection* fc) override { - return nullptr; - } + const char* getPluginVersion() const override { return "1"; } nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serial_data, size_t serial_length) override { - auto plugin = new LayerNormPluginDynamic(serial_data, serial_length); - return plugin; + return new LayerNormPluginDynamic(serial_data, serial_length); } - - void setPluginNamespace(const char* lib_namespace) override { - plugin_namespace_ = lib_namespace; - } - - const char* getPluginNamespace() const override { - return plugin_namespace_.c_str(); - } - - private: - std::string plugin_namespace_; - std::string plugin_name_; - nvinfer1::PluginFieldCollection field_collection_{0, nullptr}; - std::vector plugin_attributes_; }; REGISTER_TRT_PLUGIN_V2(LayerNormPluginDynamicCreator); From 571ca99dd0d92fc5437908575b528ea30ecb8d54 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 11:54:59 +0000 Subject: [PATCH 19/46] add pool creator and fix serialization --- .../tensorrt/plugin/pool_op_plugin.h | 36 ++++++++++++++++--- 1 file changed, 31 insertions(+), 5 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h index 26fba52aebe26..4c9f7c9e06dea 100644 --- a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h @@ -58,17 +58,16 @@ static std::vector CalcOutputSize(const std::vector& input_shape, class PoolPlugin : public PluginTensorRT { public: size_t getSerializationSize() const override { - return SerializedSize(getPluginType()) + SerializedSize(ceil_mode_) + + return getBaseSerializationSize() + SerializedSize(ceil_mode_) + SerializedSize(pool_type_) + SerializedSize(adaptive_) + SerializedSize(ksize_) + SerializedSize(strides_) + SerializedSize(paddings_) + SerializedSize(input_shape_) + - SerializedSize(output_shape_) + getBaseSerializationSize(); + SerializedSize(output_shape_); } // TRT will call this func when we need to serialize the configuration of // tensorrt. void serialize(void* buffer) const override { - SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, ceil_mode_); SerializeValue(&buffer, pool_type_); @@ -145,6 +144,20 @@ class PoolPlugin : public PluginTensorRT { std::vector output_shape_; }; +class PoolPluginCreator : public TensorRTPluginCreator { + public: + const char* getPluginName() const override { return "pool_plugin"; } + + const char* getPluginVersion() const override { return "1"; } + + nvinfer1::IPluginV2* deserializePlugin(const char* name, + const void* serial_data, + size_t serial_length) override { + return new PoolPlugin(serial_data, serial_length); + } +}; +REGISTER_TRT_PLUGIN_V2(PoolPluginCreator); + #if IS_TRT_VERSION_GE(6000) class PoolPluginDynamic : public DynamicPluginTensorRT { public: @@ -162,7 +175,6 @@ class PoolPluginDynamic : public DynamicPluginTensorRT { is_global_(is_global) {} PoolPluginDynamic(void const* serialData, size_t serialLength) { - deserializeBase(serialData, serialLength); DeserializeValue(&serialData, &serialLength, &ceil_mode_); const char* pool_type; DeserializeValue(&serialData, &serialLength, &pool_type); @@ -179,7 +191,7 @@ class PoolPluginDynamic : public DynamicPluginTensorRT { strides_, paddings_, is_global_); } - const char* getPluginType() const override { return "pool_plugin"; } + const char* getPluginType() const override { return "pool_plugin_dynamic"; } int getNbOutputs() const override { return 1; } int initialize() override { return 0; } @@ -225,6 +237,20 @@ class PoolPluginDynamic : public DynamicPluginTensorRT { std::vector paddings_; bool is_global_; }; + +class PoolPluginDynamicCreator : public TensorRTPluginCreator { + public: + const char* getPluginName() const override { return "pool_plugin_dynamic"; } + + const char* getPluginVersion() const override { return "1"; } + + nvinfer1::IPluginV2* deserializePlugin(const char* name, + const void* serial_data, + size_t serial_length) override { + return new PoolPluginDynamic(serial_data, serial_length); + } +}; +REGISTER_TRT_PLUGIN_V2(PoolPluginDynamicCreator); #endif } // namespace plugin From f60ec9cc4bc6c960778868bd2789a968b541416f Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 11:55:23 +0000 Subject: [PATCH 20/46] add prelu creator and fix serialization --- .../tensorrt/plugin/prelu_op_plugin.h | 34 ++++++++++++++++--- 1 file changed, 30 insertions(+), 4 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h index ceb4385565ff8..6dada4e687833 100644 --- a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h @@ -36,14 +36,13 @@ class PReluPlugin : public PluginTensorRT { public: size_t getSerializationSize() const override { return getBaseSerializationSize() + SerializedSize(mode_.c_str()) + - SerializedSize(weight_) + SerializedSize(getPluginType()); + SerializedSize(weight_); } // TRT will call this func when we need to serialize the configuration of // tensorrt. // It should not be called by users. void serialize(void* buffer) const override { - SerializeValue(&buffer, getPluginType()); serializeBase(buffer); SerializeValue(&buffer, weight_); SerializeValue(&buffer, mode_.c_str()); @@ -87,6 +86,20 @@ class PReluPlugin : public PluginTensorRT { void* workspace, cudaStream_t stream) override; }; +class PReluPluginCreator : public TensorRTPluginCreator { + public: + const char* getPluginName() const override { return "prelu_plugin"; } + + const char* getPluginVersion() const override { return "1"; } + + nvinfer1::IPluginV2* deserializePlugin(const char* name, + const void* serial_data, + size_t serial_length) override { + return new PReluPlugin(serial_data, serial_length); + } +}; +REGISTER_TRT_PLUGIN_V2(PReluPluginCreator); + #if IS_TRT_VERSION_GE(6000) class PReluPluginDynamic : public DynamicPluginTensorRT { public: @@ -100,7 +113,6 @@ class PReluPluginDynamic : public DynamicPluginTensorRT { // It was used for tensorrt deserialization. // It should not be called by users. PReluPluginDynamic(void const* serialData, size_t serialLength) { - deserializeBase(serialData, serialLength); DeserializeValue(&serialData, &serialLength, &weight_); const char* prelu_mode; DeserializeValue(&serialData, &serialLength, &prelu_mode); @@ -113,7 +125,7 @@ class PReluPluginDynamic : public DynamicPluginTensorRT { return ptr; } - const char* getPluginType() const override { return "prelu_plugin"; } + const char* getPluginType() const override { return "prelu_plugin_dynamic"; } int getNbOutputs() const override { return 1; } int initialize() override; void terminate() override; @@ -158,6 +170,20 @@ class PReluPluginDynamic : public DynamicPluginTensorRT { }; #endif +class PReluPluginDynamicCreator : public TensorRTPluginCreator { + public: + const char* getPluginName() const override { return "prelu_plugin_dynamic"; } + + const char* getPluginVersion() const override { return "1"; } + + nvinfer1::IPluginV2* deserializePlugin(const char* name, + const void* serial_data, + size_t serial_length) override { + return new PReluPluginDynamic(serial_data, serial_length); + } +}; +REGISTER_TRT_PLUGIN_V2(PReluPluginDynamicCreator); + } // namespace plugin } // namespace tensorrt } // namespace inference From 4c1ab091916cb8b87e81072801bab6fcb0affca4 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 11:55:47 +0000 Subject: [PATCH 21/46] add slice creator and fix serialization --- .../tensorrt/plugin/slice_op_plugin.h | 44 ++++++++----------- 1 file changed, 18 insertions(+), 26 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h index 57adbad94b01f..b656918f8fbab 100644 --- a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h @@ -66,6 +66,20 @@ class SlicePlugin : public PluginTensorRT { cudaStream_t copy_stream_; }; +class SlicePluginCreator : public TensorRTPluginCreator { + public: + const char* getPluginName() const override { return "slice_plugin"; } + + const char* getPluginVersion() const override { return "1"; } + + nvinfer1::IPluginV2* deserializePlugin(const char* name, + const void* serial_data, + size_t serial_length) override { + return new SlicePlugin(serial_data, serial_length); + } +}; +REGISTER_TRT_PLUGIN_V2(SlicePluginCreator); + #if IS_TRT_VERSION_GE(6000) class SlicePluginDynamic : public DynamicPluginTensorRT { public: @@ -78,7 +92,7 @@ class SlicePluginDynamic : public DynamicPluginTensorRT { SlicePluginDynamic(void const* serialData, size_t serialLength); - const char* getPluginType() const override { return "slice_plugin"; } + const char* getPluginType() const override { return "slice_plugin_dynamic"; } int getNbOutputs() const override { return 1; } int initialize() override; @@ -124,40 +138,18 @@ class SlicePluginDynamic : public DynamicPluginTensorRT { cudaStream_t copy_stream_; }; -class SlicePluginDynamicCreator : public nvinfer1::IPluginCreator { +class SlicePluginDynamicCreator : public TensorRTPluginCreator { public: - SlicePluginDynamicCreator() {} - const char* getPluginName() const override { return "slice_plugin"; } + const char* getPluginName() const override { return "slice_plugin_dynamic"; } const char* getPluginVersion() const override { return "1"; } - const nvinfer1::PluginFieldCollection* getFieldNames() override { - return &field_collection_; - } - - nvinfer1::IPluginV2* createPlugin( - const char* name, const nvinfer1::PluginFieldCollection* fc) override { - return nullptr; - } - nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) override { - auto plugin = new SlicePluginDynamic(serialData, serialLength); - return plugin; - } - - void setPluginNamespace(const char* libNamespace) override { - namespace_ = libNamespace; + return new SlicePluginDynamic(serialData, serialLength); } - - const char* getPluginNamespace() const override { return namespace_.c_str(); } - - private: - std::string namespace_; - nvinfer1::PluginFieldCollection field_collection_; }; - REGISTER_TRT_PLUGIN_V2(SlicePluginDynamicCreator); #endif From c7053fab1216627852363c1eb17f6ec7a079244a Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 11:56:02 +0000 Subject: [PATCH 22/46] add swish creator and fix serialization --- .../tensorrt/plugin/swish_op_plugin.h | 64 ++----------------- 1 file changed, 6 insertions(+), 58 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h index 1b8342c8c895c..8940fdce3b0b5 100644 --- a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h @@ -35,9 +35,6 @@ class SwishPlugin : public PluginTensorRT { return getBaseSerializationSize() + SerializedSize(beta_); } - // TRT will call this func when we need to serialize the configuration of - // tensorrt. - // It should not be called by users. void serialize(void* buffer) const override { serializeBase(buffer); SerializeValue(&buffer, beta_); @@ -53,7 +50,9 @@ class SwishPlugin : public PluginTensorRT { deserializeBase(serialData, serialLength); DeserializeValue(&serialData, &serialLength, &beta_); } + ~SwishPlugin() {} + int initialize() override; SwishPlugin* clone() const override { @@ -72,42 +71,17 @@ class SwishPlugin : public PluginTensorRT { void* workspace, cudaStream_t stream) override; }; -class SwishPluginCreator : public nvinfer1::IPluginCreator { +class SwishPluginCreator : public TensorRTPluginCreator { public: - SwishPluginCreator() {} const char* getPluginName() const override { return "swish_plugin"; } const char* getPluginVersion() const override { return "1"; } - const nvinfer1::PluginFieldCollection* getFieldNames() override { - return &field_collection_; - } - - nvinfer1::IPluginV2* createPlugin( - const char* name, const nvinfer1::PluginFieldCollection* fc) override { - return nullptr; - } - nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serial_data, size_t serial_length) override { - auto plugin = new SwishPlugin(serial_data, serial_length); - return plugin; - } - - void setPluginNamespace(const char* lib_namespace) override { - plugin_namespace_ = lib_namespace; + return new SwishPlugin(serial_data, serial_length); } - - const char* getPluginNamespace() const override { - return plugin_namespace_.c_str(); - } - - private: - std::string plugin_namespace_; - std::string plugin_name_; - nvinfer1::PluginFieldCollection field_collection_{0, nullptr}; - std::vector plugin_attributes_; }; REGISTER_TRT_PLUGIN_V2(SwishPluginCreator); @@ -167,44 +141,18 @@ class SwishPluginDynamic : public DynamicPluginTensorRT { float beta_; }; -class SwishPluginDynamicCreator : public nvinfer1::IPluginCreator { +class SwishPluginDynamicCreator : public TensorRTPluginCreator { public: - SwishPluginDynamicCreator() {} const char* getPluginName() const override { return "swish_plugin_dynamic"; } const char* getPluginVersion() const override { return "1"; } - const nvinfer1::PluginFieldCollection* getFieldNames() override { - return &field_collection_; - } - - nvinfer1::IPluginV2* createPlugin( - const char* name, const nvinfer1::PluginFieldCollection* fc) override { - return nullptr; - } - nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serial_data, size_t serial_length) override { - auto plugin = new SwishPluginDynamic(serial_data, serial_length); - return plugin; - } - - void setPluginNamespace(const char* lib_namespace) override { - plugin_namespace_ = lib_namespace; - } - - const char* getPluginNamespace() const override { - return plugin_namespace_.c_str(); + return new SwishPluginDynamic(serial_data, serial_length); } - - private: - std::string plugin_namespace_; - std::string plugin_name_; - nvinfer1::PluginFieldCollection field_collection_{0, nullptr}; - std::vector plugin_attributes_; }; - REGISTER_TRT_PLUGIN_V2(SwishPluginDynamicCreator); #endif From 3b2e4930f14da6e30479c0b1bfef0ca6580b841d Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 16:57:35 +0000 Subject: [PATCH 23/46] add instance norm op unittest --- .../ir/inference/test_trt_instance_norm_op.py | 93 +++++++++++++++++++ 1 file changed, 93 insertions(+) create mode 100644 python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py new file mode 100644 index 0000000000000..97cdfc7a0163f --- /dev/null +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py @@ -0,0 +1,93 @@ +# Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +from __future__ import print_function + +import unittest +import itertools +import numpy as np +from inference_pass_test import InferencePassTest +import paddle.fluid as fluid +import paddle.fluid.core as core +from paddle.fluid.core import PassVersionChecker +from paddle.fluid.core import AnalysisConfig + + +class TRTInstanceNormTest(InferencePassTest): + def setUp(self): + self.bs = 4 + self.channel = 4 + self.height = 8 + self.width = 8 + self.precision = AnalysisConfig.Precision.Float32 + self.serialize = False + self.enable_trt = True + + def build(self): + self.trt_parameters = InferencePassTest.TensorRTParam( + 1 << 30, self.bs, 3, self.precision, self.serialize, False) + + with fluid.program_guard(self.main_program, self.startup_program): + shape = [-1, self.channel, self.height, self.width] + data = fluid.data(name='in', shape=shape, dtype='float32') + instance_norm_out = fluid.layers.instance_norm(data) + out = fluid.layers.batch_norm(instance_norm_out, is_test=True) + + shape[0] = self.bs + self.feeds = {'in': np.random.random(shape).astype('float32'), } + self.fetch_list = [out] + + def check_output(self): + if core.is_compiled_with_cuda(): + use_gpu = True + atol = 1e-5 + if self.trt_parameters.precision == AnalysisConfig.Precision.Half: + atol = 2e-2 + self.check_output_with_option(use_gpu, atol, flatten=True) + self.assertTrue( + PassVersionChecker.IsCompatible('tensorrt_subgraph_pass')) + + def run_test(self): + self.build() + self.check_output() + + def run_all_tests(self): + precision_opt = [ + AnalysisConfig.Precision.Float32, AnalysisConfig.Precision.Half + ] + serialize_opt = [False, True] + + for precision, serialize in itertools.product(precision_opt, + serialize_opt): + self.precision = precision + self.serialize = serialize + self.run_test() + + def test_base(self): + self.run_test() + + def test_fp16(self): + self.precision = AnalysisConfig.Precision.Half + self.run_test() + + def test_serialize(self): + self.serialize = True + self.run_test() + + def test_all(self): + self.run_all_tests() + + +if __name__ == "__main__": + unittest.main() From 3c0a2120b8a36363b29106ba8e4933cee10058fb Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 20:04:34 +0000 Subject: [PATCH 24/46] remove redundent api --- paddle/fluid/inference/tensorrt/convert/elementwise_op.cc | 7 ++++--- paddle/fluid/inference/tensorrt/plugin/trt_plugin.h | 4 ---- 2 files changed, 4 insertions(+), 7 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc b/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc index 6820cc8cda668..45164063078b1 100644 --- a/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc +++ b/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc @@ -251,10 +251,11 @@ class ElementwiseTensorOpConverter : public OpConverter { } else { plugin::ElementWisePlugin* plugin = new plugin::ElementWisePlugin(op_type_, dims_x, dims_y, axis); - plugin->AddInput(X); - plugin->AddInput(Y); + std::vector inputs{X, Y}; + // plugin->AddInput(X); + // plugin->AddInput(Y); auto* plugin_layer = engine_->AddPlugin( - plugin->GetInputs().data(), 2, + inputs.data(), inputs.size(), reinterpret_cast(plugin)); layer = plugin_layer; diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h index 9cc99a4da3fa7..96119cae97775 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h @@ -64,10 +64,6 @@ class PluginTensorRT : public nvinfer1::IPluginV2 { nvinfer1::PluginFormat getDataFormat() const { return data_format_; } - void AddInput(nvinfer1::ITensor* input) { inputs_.push_back(input); } - - std::vector& GetInputs() { return inputs_; } - // IPluginV2 virtual const char* getPluginType() const = 0; From 3d59b633129b75ce8e32e4a82507850399456b6a Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Fri, 25 Jun 2021 20:39:35 +0000 Subject: [PATCH 25/46] fix wrong graph size to enable trt --- .../tests/unittests/ir/inference/test_trt_instance_norm_op.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py index 97cdfc7a0163f..01b686bd57bd8 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py @@ -36,7 +36,7 @@ def setUp(self): def build(self): self.trt_parameters = InferencePassTest.TensorRTParam( - 1 << 30, self.bs, 3, self.precision, self.serialize, False) + 1 << 30, self.bs, 2, self.precision, self.serialize, False) with fluid.program_guard(self.main_program, self.startup_program): shape = [-1, self.channel, self.height, self.width] From 045e90661ea6a623cdab436056c38ad87e318560 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Sat, 26 Jun 2021 08:13:27 +0000 Subject: [PATCH 26/46] instance norm function move to cc --- .../inference/tensorrt/plugin/instance_norm_op_plugin.cu | 7 +++++++ .../inference/tensorrt/plugin/instance_norm_op_plugin.h | 8 +++----- 2 files changed, 10 insertions(+), 5 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu index b9629ee7ee5ad..ef839d5d4a7bd 100644 --- a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu @@ -51,6 +51,13 @@ nvinfer1::Dims InstanceNormPlugin::getOutputDimensions( return output_dims; } +bool InstanceNormPlugin::supportsFormat(nvinfer1::DataType type, + nvinfer1::PluginFormat format) const { + return ((type == nvinfer1::DataType::kFLOAT || + type == nvinfer1::DataType::kHALF) && + (format == nvinfer1::PluginFormat::kNCHW)); +} + int InstanceNormPlugin::enqueue(int batch_size, const void *const *inputs, #if IS_TRT_VERSION_LT(8000) void **outputs, void *workspace, diff --git a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h index 16f03c7e573c7..f9dab09beebd3 100644 --- a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h @@ -88,6 +88,7 @@ class InstanceNormPlugin : public PluginTensorRT { platform::dynload::cudnnDestroyTensorDescriptor(y_desc_); platform::dynload::cudnnDestroyTensorDescriptor(b_desc_); } + int initialize() override; InstanceNormPlugin *clone() const override { @@ -98,6 +99,7 @@ class InstanceNormPlugin : public PluginTensorRT { int getNbOutputs() const override { return 1; } nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims *inputs, int nbInputDims) override; + #if IS_TRT_VERSION_LT(8000) int enqueue(int batchSize, const void *const *inputs, void **outputs, #else @@ -106,11 +108,7 @@ class InstanceNormPlugin : public PluginTensorRT { void *workspace, cudaStream_t stream) override; bool supportsFormat(nvinfer1::DataType type, - nvinfer1::PluginFormat format) const override { - return ((type == nvinfer1::DataType::kFLOAT || - type == nvinfer1::DataType::kHALF) && - (format == nvinfer1::PluginFormat::kNCHW)); - } + nvinfer1::PluginFormat format) const override; }; class InstanceNormPluginCreator : public TensorRTPluginCreator { From 3ad1eb68f72ec0803ba714f18fda14d10e049ac3 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Sat, 26 Jun 2021 08:20:19 +0000 Subject: [PATCH 27/46] add trt elementwise ut to trigger coverage --- .../ir/inference/test_trt_elementwise_op.py | 56 +++++++++++++++++++ 1 file changed, 56 insertions(+) create mode 100644 python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py new file mode 100644 index 0000000000000..02f44b4fcb116 --- /dev/null +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py @@ -0,0 +1,56 @@ +# Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +from __future__ import print_function + +import unittest +import numpy as np +from inference_pass_test import InferencePassTest +import paddle.fluid as fluid +import paddle.fluid.core as core +from paddle.fluid.core import PassVersionChecker +from paddle.fluid.core import AnalysisConfig + + +class TensorRTSubgraphPassElementwiseBroadcastTest(InferencePassTest): + def setUp(self): + with fluid.program_guard(self.main_program, self.startup_program): + data1 = fluid.data( + name="data1", shape=[-1, 3, 64, 64], dtype="float32") + data2 = fluid.data( + name="data2", shape=[-1, 3, 64, 1], dtype="float32") + eltwise_out = self.append_eltwise(data1, data2) + out = fluid.layers.batch_norm(eltwise_out, is_test=True) + self.feeds = { + "data1": np.random.random([1, 3, 64, 64]).astype("float32"), + "data2": np.random.random([1, 3, 64, 1]).astype("float32"), + } + self.enable_trt = True + self.trt_parameters = TensorRTSubgraphPassElementwiseBroadcastTest.TensorRTParam( + 1 << 30, 32, 0, AnalysisConfig.Precision.Float32, True, False) + self.fetch_list = [out] + + def append_eltwise(self, data1, data2): + return fluid.layers.elementwise_add(x=data1, y=data2, axis=0) + + def test_check_output(self): + if core.is_compiled_with_cuda(): + use_gpu = True + self.check_output_with_option(use_gpu) + self.assertTrue( + PassVersionChecker.IsCompatible('tensorrt_subgraph_pass')) + + +if __name__ == "__main__": + unittest.main() From 92849b96dbee3f572b3464ddacb89ea6b51e2354 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Sat, 26 Jun 2021 10:35:36 +0000 Subject: [PATCH 28/46] remove opt cahce to hit serialization coverage --- .../tests/unittests/ir/inference/test_trt_elementwise_op.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py index 02f44b4fcb116..f84202df5fb93 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py @@ -14,6 +14,8 @@ from __future__ import print_function +import os +import shutil import unittest import numpy as np from inference_pass_test import InferencePassTest @@ -45,6 +47,8 @@ def append_eltwise(self, data1, data2): return fluid.layers.elementwise_add(x=data1, y=data2, axis=0) def test_check_output(self): + if os.path.exists(self.path + "_opt_cache"): + shutil.rmtree(self.path + "_opt_cache") if core.is_compiled_with_cuda(): use_gpu = True self.check_output_with_option(use_gpu) From 3723bbb0d3ecb8ff2bf7cd8c01999d534dd73a8e Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Sat, 26 Jun 2021 10:35:47 +0000 Subject: [PATCH 29/46] remove opt cahce to hit serialization coverage --- .../ir/inference/test_trt_instance_norm_op.py | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py index 01b686bd57bd8..d283465dcba09 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py @@ -14,6 +14,8 @@ from __future__ import print_function +import os +import shutil import unittest import itertools import numpy as np @@ -48,7 +50,9 @@ def build(self): self.feeds = {'in': np.random.random(shape).astype('float32'), } self.fetch_list = [out] - def check_output(self): + def check_output(self, remove_cache=False): + if remove_cache and os.path.exists(self.path + "_opt_cache"): + shutil.rmtree(self.path + "_opt_cache") if core.is_compiled_with_cuda(): use_gpu = True atol = 1e-5 @@ -58,9 +62,9 @@ def check_output(self): self.assertTrue( PassVersionChecker.IsCompatible('tensorrt_subgraph_pass')) - def run_test(self): + def run_test(self, remove_cache=False): self.build() - self.check_output() + self.check_output(remove_cache) def run_all_tests(self): precision_opt = [ @@ -83,7 +87,7 @@ def test_fp16(self): def test_serialize(self): self.serialize = True - self.run_test() + self.run_test(remove_cache=True) def test_all(self): self.run_all_tests() From 5a6b3298cf654a324ae3ddfeb618f181c676ec13 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Sat, 26 Jun 2021 10:55:13 +0000 Subject: [PATCH 30/46] remove unused code --- .../inference/tensorrt/plugin/trt_plugin.cc | 59 ++++--------------- .../inference/tensorrt/plugin/trt_plugin.h | 27 +-------- 2 files changed, 14 insertions(+), 72 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc index 64cf94fc93357..7288797a9bc3c 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc @@ -21,10 +21,9 @@ namespace plugin { inline void Seria(void*& buffer, // NOLINT const std::vector& input_dims, - size_t max_batch_size, nvinfer1::DataType data_type, + nvinfer1::DataType data_type, nvinfer1::PluginFormat data_format, bool with_fp16) { SerializeValue(&buffer, input_dims); - SerializeValue(&buffer, max_batch_size); SerializeValue(&buffer, data_type); SerializeValue(&buffer, data_format); SerializeValue(&buffer, with_fp16); @@ -32,37 +31,33 @@ inline void Seria(void*& buffer, // NOLINT inline void Deseria(void const*& serial_data, size_t& serial_length, // NOLINT std::vector* input_dims, - size_t* max_batch_size, nvinfer1::DataType* data_type, + nvinfer1::DataType* data_type, nvinfer1::PluginFormat* data_format, bool* with_fp16) { DeserializeValue(&serial_data, &serial_length, input_dims); - DeserializeValue(&serial_data, &serial_length, max_batch_size); DeserializeValue(&serial_data, &serial_length, data_type); DeserializeValue(&serial_data, &serial_length, data_format); DeserializeValue(&serial_data, &serial_length, with_fp16); } inline size_t SeriaSize(const std::vector& input_dims, - size_t max_batch_size, nvinfer1::DataType data_type, + nvinfer1::DataType data_type, nvinfer1::PluginFormat data_format, bool with_fp16) { - return (SerializedSize(input_dims) + SerializedSize(max_batch_size) + - SerializedSize(data_type) + SerializedSize(data_format) + - SerializedSize(with_fp16)); + return (SerializedSize(input_dims) + SerializedSize(data_type) + + SerializedSize(data_format) + SerializedSize(with_fp16)); } void PluginTensorRT::serializeBase(void*& buffer) const { - Seria(buffer, input_dims_, max_batch_size_, data_type_, data_format_, - with_fp16_); + Seria(buffer, input_dims_, data_type_, data_format_, with_fp16_); } void PluginTensorRT::deserializeBase(void const*& serial_data, size_t& serial_length) { - Deseria(serial_data, serial_length, &input_dims_, &max_batch_size_, - &data_type_, &data_format_, &with_fp16_); + Deseria(serial_data, serial_length, &input_dims_, &data_type_, &data_format_, + &with_fp16_); } size_t PluginTensorRT::getBaseSerializationSize() const { - return SeriaSize(input_dims_, max_batch_size_, data_type_, data_format_, - with_fp16_); + return SeriaSize(input_dims_, data_type_, data_format_, with_fp16_); } bool PluginTensorRT::supportsFormat(nvinfer1::DataType type, @@ -78,47 +73,20 @@ void PluginTensorRT::configureWithFormat( data_type_ = type; data_format_ = format; input_dims_.assign(input_dims, input_dims + num_inputs); - max_batch_size_ = max_batch_size; -} -/* -nvinfer1::DataType PluginTensorRT::getOutputDataType( - int32_t index, const nvinfer1::DataType* input_types, - int32_t nb_inputs) const { - return input_types[0]; -} - -bool PluginTensorRT::isOutputBroadcastAcrossBatch( - int32_t output_index, const bool* input_is_broadcasted, - int32_t nb_inputs) const { - return false; -} - -bool PluginTensorRT::canBroadcastInputAcrossBatch(int32_t input_index) const { - return false; } -void PluginTensorRT::configurePlugin( - const nvinfer1::Dims* input_dims, int32_t nb_inputs, - const nvinfer1::Dims* output_dims, int32_t nb_outputs, - const nvinfer1::DataType* input_types, - const nvinfer1::DataType* output_types, const bool* input_is_broadcast, - const bool* output_is_broadcast, nvinfer1::PluginFormat float_format, - int32_t max_batch_size) {} -*/ void PluginTensorRTV2Ext::serializeBase(void*& buffer) const { - Seria(buffer, input_dims_, max_batch_size_, data_type_, data_format_, - with_fp16_); + Seria(buffer, input_dims_, data_type_, data_format_, with_fp16_); } void PluginTensorRTV2Ext::deserializeBase(void const*& serial_data, size_t& serial_length) { - Deseria(serial_data, serial_length, &input_dims_, &max_batch_size_, - &data_type_, &data_format_, &with_fp16_); + Deseria(serial_data, serial_length, &input_dims_, &data_type_, &data_format_, + &with_fp16_); } size_t PluginTensorRTV2Ext::getBaseSerializationSize() const { - return SeriaSize(input_dims_, max_batch_size_, data_type_, data_format_, - with_fp16_); + return SeriaSize(input_dims_, data_type_, data_format_, with_fp16_); } void PluginTensorRTV2Ext::configurePlugin( @@ -129,7 +97,6 @@ void PluginTensorRTV2Ext::configurePlugin( const bool* output_is_broadcast, nvinfer1::PluginFormat float_format, int32_t max_batch_size) { input_dims_.assign(input_dims, input_dims + nb_inputs); - max_batch_size_ = max_batch_size; data_format_ = float_format; data_type_ = input_types[0]; } diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h index 96119cae97775..40458ab7d925b 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h @@ -58,8 +58,6 @@ class PluginTensorRT : public nvinfer1::IPluginV2 { return input_dims_.at(index); } - size_t getMaxBatchSize() const { return max_batch_size_; } - nvinfer1::DataType getDataType() const { return data_type_; } nvinfer1::PluginFormat getDataFormat() const { return data_format_; } @@ -121,27 +119,7 @@ class PluginTensorRT : public nvinfer1::IPluginV2 { } const char* getPluginNamespace() const override { return namespace_.c_str(); } - /* - // IPluginV2Ext - nvinfer1::DataType getOutputDataType(int32_t index, - const nvinfer1::DataType* input_types, - int32_t nb_inputs) const override; - - bool isOutputBroadcastAcrossBatch(int32_t output_index, - const bool* input_is_broadcasted, - int32_t nb_inputs) const override; - - bool canBroadcastInputAcrossBatch(int32_t input_index) const override; - - void configurePlugin(const nvinfer1::Dims* input_dims, int32_t nb_inputs, - const nvinfer1::Dims* output_dims, int32_t nb_outputs, - const nvinfer1::DataType* input_types, - const nvinfer1::DataType* output_types, - const bool* input_is_broadcast, - const bool* output_is_broadcast, - nvinfer1::PluginFormat float_format, - int32_t max_batch_size) override; - */ + protected: // Deserialize input_dims, max_batch_size, data_type, data_format void deserializeBase(void const*& serial_data, // NOLINT @@ -151,7 +129,6 @@ class PluginTensorRT : public nvinfer1::IPluginV2 { void serializeBase(void*& buffer) const; // NOLINT std::vector input_dims_; - size_t max_batch_size_; nvinfer1::DataType data_type_; nvinfer1::PluginFormat data_format_; @@ -172,7 +149,6 @@ class PluginTensorRTV2Ext : public nvinfer1::IPluginV2Ext { nvinfer1::Dims const& getInputDims(int index) const { return input_dims_.at(index); } - size_t getMaxBatchSize() const { return max_batch_size_; } nvinfer1::DataType getDataType() const { return data_type_; } nvinfer1::PluginFormat getDataFormat() const { return data_format_; } @@ -265,7 +241,6 @@ class PluginTensorRTV2Ext : public nvinfer1::IPluginV2Ext { protected: std::vector input_dims_; - size_t max_batch_size_; nvinfer1::DataType data_type_; nvinfer1::PluginFormat data_format_; std::vector inputs_; From b57de2525702a1cf4533f7372c9784ecbd6eff02 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Sat, 26 Jun 2021 13:11:42 +0000 Subject: [PATCH 31/46] remove unused inputs_ --- paddle/fluid/inference/tensorrt/plugin/trt_plugin.h | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h index 40458ab7d925b..a24a5bdf8e963 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h @@ -48,6 +48,7 @@ typedef std::function PluginConstructFunc; class PluginTensorRT : public nvinfer1::IPluginV2 { public: PluginTensorRT() : with_fp16_(false) {} + // It was used for TensorRT deserialization. // It should not be called by users. PluginTensorRT(const void* serialized_data, size_t length) {} @@ -73,7 +74,7 @@ class PluginTensorRT : public nvinfer1::IPluginV2 { const nvinfer1::Dims* input_dims, int num_inputs) = 0; - // Check format support. The default is FLOAT32 and NCHW. + // Check format support. The default is FLOAT32 and kLINEAR. bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const override; @@ -132,7 +133,6 @@ class PluginTensorRT : public nvinfer1::IPluginV2 { nvinfer1::DataType data_type_; nvinfer1::PluginFormat data_format_; - std::vector inputs_; bool with_fp16_; private: @@ -243,7 +243,6 @@ class PluginTensorRTV2Ext : public nvinfer1::IPluginV2Ext { std::vector input_dims_; nvinfer1::DataType data_type_; nvinfer1::PluginFormat data_format_; - std::vector inputs_; bool with_fp16_; private: From 4a03a9cd68d37941d63dc9da7b2a688e27dbf41f Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Sat, 26 Jun 2021 13:25:50 +0000 Subject: [PATCH 32/46] add dbg info --- paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h | 3 +++ 1 file changed, 3 insertions(+) diff --git a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h index 5dd3142c75839..11e690e49be96 100644 --- a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h @@ -14,6 +14,7 @@ limitations under the License. */ #pragma once +#include #include #include #include "paddle/fluid/inference/tensorrt/plugin/trt_plugin.h" @@ -68,6 +69,7 @@ class ElementWisePlugin : public PluginTensorRT { void* workspace, cudaStream_t stream); size_t getSerializationSize() const override { + VLOG(2) << "ElementWisePlugin getSerializationSize"; return getBaseSerializationSize() + SerializedSize(type_.c_str()) + SerializedSize(dims_x_) + SerializedSize(dims_y_) + SerializedSize(axis_) + SerializedSize(prev_size_) + @@ -75,6 +77,7 @@ class ElementWisePlugin : public PluginTensorRT { } void serialize(void* buffer) const override { + VLOG(2) << "ElementWisePlugin serialize"; serializeBase(buffer); SerializeValue(&buffer, type_.c_str()); SerializeValue(&buffer, dims_x_); From 1d5c960ef4abb8f19bb669a047fac3916e16d2ae Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Sat, 26 Jun 2021 15:24:17 +0000 Subject: [PATCH 33/46] remove dbg info --- paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h | 2 -- 1 file changed, 2 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h index 11e690e49be96..836f09bb00103 100644 --- a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h @@ -69,7 +69,6 @@ class ElementWisePlugin : public PluginTensorRT { void* workspace, cudaStream_t stream); size_t getSerializationSize() const override { - VLOG(2) << "ElementWisePlugin getSerializationSize"; return getBaseSerializationSize() + SerializedSize(type_.c_str()) + SerializedSize(dims_x_) + SerializedSize(dims_y_) + SerializedSize(axis_) + SerializedSize(prev_size_) + @@ -77,7 +76,6 @@ class ElementWisePlugin : public PluginTensorRT { } void serialize(void* buffer) const override { - VLOG(2) << "ElementWisePlugin serialize"; serializeBase(buffer); SerializeValue(&buffer, type_.c_str()); SerializeValue(&buffer, dims_x_); From d29a365b53fb567a12351aae529f887e22e82a3b Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Sat, 26 Jun 2021 15:25:18 +0000 Subject: [PATCH 34/46] add instance norm serialization --- .../inference/tests/api/trt_instance_norm_converter_test.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/paddle/fluid/inference/tests/api/trt_instance_norm_converter_test.cc b/paddle/fluid/inference/tests/api/trt_instance_norm_converter_test.cc index ceb8b99774e48..06ac50fb9ecc5 100644 --- a/paddle/fluid/inference/tests/api/trt_instance_norm_converter_test.cc +++ b/paddle/fluid/inference/tests/api/trt_instance_norm_converter_test.cc @@ -29,7 +29,7 @@ TEST(TensorRT, instance_norm) { config.SetModel(model_dir); config.SwitchUseFeedFetchOps(false); config.EnableTensorRtEngine(1 << 20, batch_size, 0, - AnalysisConfig::Precision::kFloat32, false); + AnalysisConfig::Precision::kFloat32, true); auto predictor = CreatePaddlePredictor(config); From 6ac45e290edaaca1d220daba784a157664e8bf50 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Sat, 26 Jun 2021 17:18:53 +0000 Subject: [PATCH 35/46] roll back --- paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h | 1 - .../inference/tests/api/trt_instance_norm_converter_test.cc | 2 +- 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h index 836f09bb00103..5dd3142c75839 100644 --- a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h @@ -14,7 +14,6 @@ limitations under the License. */ #pragma once -#include #include #include #include "paddle/fluid/inference/tensorrt/plugin/trt_plugin.h" diff --git a/paddle/fluid/inference/tests/api/trt_instance_norm_converter_test.cc b/paddle/fluid/inference/tests/api/trt_instance_norm_converter_test.cc index 06ac50fb9ecc5..ceb8b99774e48 100644 --- a/paddle/fluid/inference/tests/api/trt_instance_norm_converter_test.cc +++ b/paddle/fluid/inference/tests/api/trt_instance_norm_converter_test.cc @@ -29,7 +29,7 @@ TEST(TensorRT, instance_norm) { config.SetModel(model_dir); config.SwitchUseFeedFetchOps(false); config.EnableTensorRtEngine(1 << 20, batch_size, 0, - AnalysisConfig::Precision::kFloat32, true); + AnalysisConfig::Precision::kFloat32, false); auto predictor = CreatePaddlePredictor(config); From 4397920167b3cf078d3332a38eb87aa2ad8849b2 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Mon, 28 Jun 2021 08:05:17 +0000 Subject: [PATCH 36/46] remove comment code --- paddle/fluid/inference/tensorrt/convert/elementwise_op.cc | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc b/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc index 45164063078b1..2f802ea8d181e 100644 --- a/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc +++ b/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc @@ -251,9 +251,8 @@ class ElementwiseTensorOpConverter : public OpConverter { } else { plugin::ElementWisePlugin* plugin = new plugin::ElementWisePlugin(op_type_, dims_x, dims_y, axis); + std::vector inputs{X, Y}; - // plugin->AddInput(X); - // plugin->AddInput(Y); auto* plugin_layer = engine_->AddPlugin( inputs.data(), inputs.size(), reinterpret_cast(plugin)); From faa9bf60c34a95a28cc91781e62c626302419ce0 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Mon, 28 Jun 2021 08:37:39 +0000 Subject: [PATCH 37/46] remove trt plugin registery --- paddle/fluid/inference/tensorrt/engine.h | 2 +- .../inference/tensorrt/plugin/CMakeLists.txt | 2 +- .../plugin/anchor_generator_op_plugin.cu | 3 +- .../tensorrt/plugin/elementwise_op_plugin.cu | 2 +- .../plugin/emb_eltwise_layernorm_plugin.cu | 2 +- .../tensorrt/plugin/gather_nd_op_plugin.cu | 2 +- .../tensorrt/plugin/gelu_op_plugin.cu | 2 +- .../tensorrt/plugin/hard_swish_op_plugin.cu | 2 +- .../plugin/instance_norm_op_plugin.cu | 2 +- .../tensorrt/plugin/layer_norm_op_plugin.cu | 2 +- .../tensorrt/plugin/pool_op_plugin.cu | 2 +- .../tensorrt/plugin/prelu_op_plugin.cu | 2 +- .../tensorrt/plugin/qkv_to_context_plugin.cu | 2 +- .../tensorrt/plugin/roi_align_op_plugin.cu | 2 +- .../plugin/skip_layernorm_op_plugin.cu | 2 +- .../tensorrt/plugin/slice_op_plugin.cu | 2 +- .../tensorrt/plugin/special_slice_plugin.cu | 2 +- .../tensorrt/plugin/split_op_plugin.cu | 2 +- .../tensorrt/plugin/stack_op_plugin.cu | 2 +- .../tensorrt/plugin/swish_op_plugin.cu | 2 +- .../tensorrt/plugin/trt_plugin_factory.cc | 54 ------------- .../tensorrt/plugin/trt_plugin_factory.h | 81 ------------------- .../tensorrt/plugin/yolo_box_op_plugin.cu | 2 +- 23 files changed, 21 insertions(+), 157 deletions(-) delete mode 100644 paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.cc delete mode 100644 paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h diff --git a/paddle/fluid/inference/tensorrt/engine.h b/paddle/fluid/inference/tensorrt/engine.h index f2fbacf151a52..5c574f2b57602 100644 --- a/paddle/fluid/inference/tensorrt/engine.h +++ b/paddle/fluid/inference/tensorrt/engine.h @@ -1,4 +1,5 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. + Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. @@ -30,7 +31,6 @@ limitations under the License. */ #include "paddle/fluid/inference/engine.h" #include "paddle/fluid/inference/tensorrt/helper.h" #include "paddle/fluid/inference/tensorrt/plugin/trt_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" #include "paddle/fluid/inference/tensorrt/trt_int8_calibrator.h" #include "paddle/fluid/inference/utils/singleton.h" diff --git a/paddle/fluid/inference/tensorrt/plugin/CMakeLists.txt b/paddle/fluid/inference/tensorrt/plugin/CMakeLists.txt index 26125d21ad7d1..311c2312a9f45 100644 --- a/paddle/fluid/inference/tensorrt/plugin/CMakeLists.txt +++ b/paddle/fluid/inference/tensorrt/plugin/CMakeLists.txt @@ -1,6 +1,6 @@ nv_library(tensorrt_plugin SRCS trt_plugin.cc split_op_plugin.cu elementwise_op_plugin.cu - prelu_op_plugin.cu trt_plugin_factory.cc gelu_op_plugin.cu + prelu_op_plugin.cu gelu_op_plugin.cu pool_op_plugin.cu swish_op_plugin.cu layer_norm_op_plugin.cu instance_norm_op_plugin.cu emb_eltwise_layernorm_plugin.cu qkv_to_context_plugin.cu skip_layernorm_op_plugin.cu slice_op_plugin.cu diff --git a/paddle/fluid/inference/tensorrt/plugin/anchor_generator_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/anchor_generator_op_plugin.cu index 8e9845183b3fe..cadfa1185c5be 100644 --- a/paddle/fluid/inference/tensorrt/plugin/anchor_generator_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/anchor_generator_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -18,8 +19,6 @@ #include #include "paddle/fluid/inference/tensorrt/plugin/anchor_generator_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" - #include "paddle/fluid/operators/detection/anchor_generator_op.h" namespace paddle { diff --git a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu index 274adb535fc69..5bbab49eaf558 100644 --- a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu @@ -1,4 +1,5 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. + Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. @@ -14,7 +15,6 @@ limitations under the License. */ #include #include "paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" namespace paddle { namespace inference { diff --git a/paddle/fluid/inference/tensorrt/plugin/emb_eltwise_layernorm_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/emb_eltwise_layernorm_plugin.cu index c873b1fc310de..9105619660189 100644 --- a/paddle/fluid/inference/tensorrt/plugin/emb_eltwise_layernorm_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/emb_eltwise_layernorm_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -20,7 +21,6 @@ #include "paddle/fluid/framework/tensor.h" #include "paddle/fluid/framework/tensor_util.h" #include "paddle/fluid/inference/tensorrt/plugin/emb_eltwise_layernorm_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" #include "paddle/fluid/operators/math/bert_encoder_functor.h" namespace paddle { diff --git a/paddle/fluid/inference/tensorrt/plugin/gather_nd_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/gather_nd_op_plugin.cu index 5f4ac054c95b3..c9d30f7ccf547 100644 --- a/paddle/fluid/inference/tensorrt/plugin/gather_nd_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/gather_nd_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2021 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -22,7 +23,6 @@ #include "NvInferRuntimeCommon.h" #include "paddle/fluid/inference/tensorrt/plugin/gather_nd_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" #include "paddle/fluid/platform/place.h" namespace paddle { diff --git a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu index a301fa37b2fa7..2275c742d6d28 100644 --- a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -16,7 +17,6 @@ #include #include #include "paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" #include "paddle/fluid/platform/float16.h" namespace paddle { diff --git a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu index 139e650aac961..3e37f04370d76 100644 --- a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -15,7 +16,6 @@ #include #include #include "paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" namespace paddle { namespace inference { diff --git a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu index a5d866d1d9cf2..4d7720d69d1a7 100644 --- a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -17,7 +18,6 @@ #include #include "glog/logging.h" #include "paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" #include "paddle/fluid/platform/cudnn_helper.h" namespace paddle { diff --git a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu index 64a4c553d1a94..a2fbfe1f7d74f 100644 --- a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -17,7 +18,6 @@ #include #include "glog/logging.h" #include "paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" #include "paddle/fluid/operators/layer_norm_op.h" namespace paddle { diff --git a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu index be577a52e2d89..85fb495d64162 100644 --- a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -13,7 +14,6 @@ // limitations under the License. #include "paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" #include "paddle/fluid/operators/math/pooling.h" namespace paddle { diff --git a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu index 93c9ef3a77914..7082e685ac009 100644 --- a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -19,7 +20,6 @@ #include "glog/logging.h" #include "paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" #include "paddle/fluid/operators/math/prelu.h" namespace paddle { diff --git a/paddle/fluid/inference/tensorrt/plugin/qkv_to_context_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/qkv_to_context_plugin.cu index 5f10e5821c4f7..c9ae888a170a0 100644 --- a/paddle/fluid/inference/tensorrt/plugin/qkv_to_context_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/qkv_to_context_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -20,7 +21,6 @@ #include "paddle/fluid/framework/tensor.h" #include "paddle/fluid/framework/tensor_util.h" #include "paddle/fluid/inference/tensorrt/plugin/qkv_to_context_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" #include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_utils.h" #include "paddle/fluid/operators/math/bert_encoder_functor.h" #include "paddle/fluid/operators/math/blas.h" diff --git a/paddle/fluid/inference/tensorrt/plugin/roi_align_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/roi_align_op_plugin.cu index 6e7ed0054f502..39feded7e1ae6 100644 --- a/paddle/fluid/inference/tensorrt/plugin/roi_align_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/roi_align_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -17,7 +18,6 @@ #include #include "paddle/fluid/inference/tensorrt/plugin/roi_align_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" namespace paddle { namespace inference { diff --git a/paddle/fluid/inference/tensorrt/plugin/skip_layernorm_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/skip_layernorm_op_plugin.cu index 7be9e3a740ab1..6a8922f6b6b97 100644 --- a/paddle/fluid/inference/tensorrt/plugin/skip_layernorm_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/skip_layernorm_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -19,7 +20,6 @@ #include #include "glog/logging.h" #include "paddle/fluid/inference/tensorrt/plugin/skip_layernorm_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" #include "paddle/fluid/operators/math/bert_encoder_functor.h" namespace paddle { diff --git a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu index c38391944264e..862ddc009eb4b 100644 --- a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -19,7 +20,6 @@ #include #include "glog/logging.h" #include "paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" namespace paddle { namespace inference { diff --git a/paddle/fluid/inference/tensorrt/plugin/special_slice_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/special_slice_plugin.cu index fdb14f9ceaf29..b7c76335ccb8e 100644 --- a/paddle/fluid/inference/tensorrt/plugin/special_slice_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/special_slice_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -16,7 +17,6 @@ #include #include #include "paddle/fluid/inference/tensorrt/plugin/special_slice_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" namespace paddle { namespace inference { diff --git a/paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu index 24d4715e0312d..85ae595e7588f 100644 --- a/paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -15,7 +16,6 @@ #include #include #include "paddle/fluid/inference/tensorrt/plugin/split_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" namespace paddle { namespace inference { diff --git a/paddle/fluid/inference/tensorrt/plugin/stack_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/stack_op_plugin.cu index 79ec2066faa13..c796990d9ef1d 100644 --- a/paddle/fluid/inference/tensorrt/plugin/stack_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/stack_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -16,7 +17,6 @@ #include #include #include "paddle/fluid/inference/tensorrt/plugin/stack_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" namespace paddle { namespace inference { diff --git a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu index 4af08cda5516f..70b2e48c23d21 100644 --- a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -17,7 +18,6 @@ #include #include "glog/logging.h" #include "paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" namespace paddle { namespace inference { diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.cc b/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.cc deleted file mode 100644 index 94f5beccd56cb..0000000000000 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.cc +++ /dev/null @@ -1,54 +0,0 @@ -// Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. - -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" - -#include "paddle/fluid/platform/enforce.h" - -namespace paddle { -namespace inference { -namespace tensorrt { -namespace plugin { - -#if false -PluginTensorRT* PluginFactoryTensorRT::createPlugin(const char* layer_name, - const void* serial_data, - size_t serial_length) { - const char* plugin_type; - DeserializeValue(&serial_data, &serial_length, &plugin_type); - - PADDLE_ENFORCE_EQ( - Has(plugin_type), true, - platform::errors::NotFound("TensorRT plugin type `%s` does not exists.", - plugin_type)); - auto plugin = plugin_registry_[plugin_type](serial_data, serial_length); - owned_plugins_.emplace_back(plugin); - - return plugin; -} - -bool PluginFactoryTensorRT::RegisterPlugin( - const std::string& op_name, PluginDeserializeFunc deserialize_func) { - if (Has(op_name)) return false; - auto ret = plugin_registry_.emplace(op_name, deserialize_func); - return ret.second; -} - -void PluginFactoryTensorRT::DestroyPlugins() { owned_plugins_.clear(); } -#endif - -} // namespace plugin -} // namespace tensorrt -} // namespace inference -} // namespace paddle diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h b/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h deleted file mode 100644 index 98ea916b870c9..0000000000000 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h +++ /dev/null @@ -1,81 +0,0 @@ -// Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. - -#pragma once - -#include -#include -#include -#include -#include -#include -#include - -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin.h" -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_utils.h" -#include "paddle/fluid/inference/utils/singleton.h" -#include "paddle/fluid/platform/enforce.h" -#include "paddle/fluid/platform/variant.h" - -namespace paddle { -namespace inference { -namespace tensorrt { -namespace plugin { - -#if false -class PluginFactoryTensorRT : public nvinfer1::IPluginFactory, - public DeleteHelper { - public: - // Deserialization method - PluginTensorRT* createPlugin(const char* layer_name, const void* serial_data, - size_t serial_length) override; - - bool RegisterPlugin(const std::string& op_name, - PluginDeserializeFunc deserialize_func); - - bool Has(const std::string& op_name) { - return plugin_registry_.find(op_name) != plugin_registry_.end(); - } - - void DestroyPlugins(); - - protected: - std::unordered_map plugin_registry_; - - std::list> owned_plugins_; -}; - -class TrtPluginRegistrar { - public: - TrtPluginRegistrar(const std::string& name, - PluginDeserializeFunc deserialize_func) { - inference::Singleton::Global().RegisterPlugin( - name, deserialize_func); - } -}; - -#define REGISTER_TRT_PLUGIN(name, deserialize_func) \ - REGISTER_TRT_PLUGIN_UNIQ(__COUNTER__, name, deserialize_func) - -#define REGISTER_TRT_PLUGIN_UNIQ(ctr, name, deserialize_func) \ - static paddle::inference::tensorrt::plugin::TrtPluginRegistrar \ - trt_plugin_registrar##ctr UNUSED = \ - paddle::inference::tensorrt::plugin::TrtPluginRegistrar( \ - name, deserialize_func) -#endif - -} // namespace plugin -} // namespace tensorrt -} // namespace inference -} // namespace paddle diff --git a/paddle/fluid/inference/tensorrt/plugin/yolo_box_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/yolo_box_op_plugin.cu index f9767f3855948..de635a7f08530 100644 --- a/paddle/fluid/inference/tensorrt/plugin/yolo_box_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/yolo_box_op_plugin.cu @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -17,7 +18,6 @@ #include #include -#include "paddle/fluid/inference/tensorrt/plugin/trt_plugin_factory.h" #include "paddle/fluid/inference/tensorrt/plugin/yolo_box_op_plugin.h" #include "paddle/fluid/operators/detection/yolo_box_op.h" From 8efb2c5fefa4f426a5ee248d58a7c821f5eaac58 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Mon, 28 Jun 2021 10:07:30 +0000 Subject: [PATCH 38/46] fix prelu dynamic serialization --- .../tensorrt/plugin/prelu_op_plugin.cu | 18 ++++++++++++++++-- .../tensorrt/plugin/prelu_op_plugin.h | 10 ++-------- 2 files changed, 18 insertions(+), 10 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu index 7082e685ac009..a707f7c030ce5 100644 --- a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu @@ -99,9 +99,23 @@ int PReluPluginDynamic::initialize() { cudaMemcpyHostToDevice); return 0; } -size_t PReluPluginDynamic::getSerializationSize() const { return 0; } -void PReluPluginDynamic::serialize(void *buffer) const {} +PReluPluginDynamic::PReluPluginDynamic(void const *serialData, + size_t serialLength) { + DeserializeValue(&serialData, &serialLength, &weight_); + const char *prelu_mode; + DeserializeValue(&serialData, &serialLength, &prelu_mode); + mode_ = std::string(prelu_mode); +} + +size_t PReluPluginDynamic::getSerializationSize() const { + return SerializedSize(mode_.c_str()) + SerializedSize(weight_); +} + +void PReluPluginDynamic::serialize(void *buffer) const { + SerializeValue(&buffer, weight_); + SerializeValue(&buffer, mode_.c_str()); +} nvinfer1::DimsExprs PReluPluginDynamic::getOutputDimensions( int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs, diff --git a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h index 6dada4e687833..3bc126f75224b 100644 --- a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -110,14 +111,7 @@ class PReluPluginDynamic : public DynamicPluginTensorRT { std::copy(weight, weight + weight_num, weight_.data()); } - // It was used for tensorrt deserialization. - // It should not be called by users. - PReluPluginDynamic(void const* serialData, size_t serialLength) { - DeserializeValue(&serialData, &serialLength, &weight_); - const char* prelu_mode; - DeserializeValue(&serialData, &serialLength, &prelu_mode); - mode_ = std::string(prelu_mode); - } + PReluPluginDynamic(void const* serialData, size_t serialLength); ~PReluPluginDynamic() {} nvinfer1::IPluginV2DynamicExt* clone() const override { auto ptr = new PReluPluginDynamic(weight_.data(), weight_.size(), mode_); From 55fb335251f0bf5ef7a64b0a50af1284eb0bb56f Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Mon, 28 Jun 2021 10:08:28 +0000 Subject: [PATCH 39/46] add prelu ut and reduce the input size to reduce memory usage --- .../ir/inference/test_trt_activation_pass.py | 69 ++++++++++++++++++- 1 file changed, 67 insertions(+), 2 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_activation_pass.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_activation_pass.py index f71951497f2af..8e196f5081f73 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_activation_pass.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_activation_pass.py @@ -33,11 +33,11 @@ def setUp(self): self.setUpTensorRTParam() with fluid.program_guard(self.main_program, self.startup_program): data = fluid.data( - name="data", shape=[-1, 6, 64, 64], dtype="float32") + name="data", shape=[-1, 6, 32, 32], dtype="float32") act_out = self.append_act(data) out = fluid.layers.batch_norm(act_out, is_test=True) self.feeds = { - "data": np.random.random([1, 6, 64, 64]).astype("float32"), + "data": np.random.random([1, 6, 32, 32]).astype("float32"), } self.fetch_list = [out] @@ -154,6 +154,71 @@ def append_act(self, x): return fluid.layers.prelu(x, mode='element') +class TensorRTSubgraphPassPreluDynamicTest(TensorRTSubgraphPassActivationTest): + def setUpTensorRTParam(self): + self.enable_trt = True + self.trt_parameters = TensorRTSubgraphPassActivationTest.TensorRTParam( + 1 << 30, 32, 0, AnalysisConfig.Precision.Float32, False, False) + self.dynamic_shape_params = TensorRTSubgraphPassActivationTest.DynamicShapeParam( + { + 'data': [1, 6, 8, 8] + }, {'data': [1, 6, 512, 512]}, {'data': [1, 6, 256, 256]}, False) + + def append_act(self, x): + return fluid.layers.prelu(x, mode='all') + + +class TensorRTSubgraphPassPreluFp16Test(TensorRTSubgraphPassActivationTest): + def setUpTensorRTParam(self): + self.enable_trt = True + self.trt_parameters = TensorRTSubgraphPassActivationTest.TensorRTParam( + 1 << 30, 32, 0, AnalysisConfig.Precision.Half, False, False) + + def append_act(self, x): + return fluid.layers.prelu(x, mode='all') + + +class TensorRTSubgraphPassPreluFp16SerializeTest( + TensorRTSubgraphPassActivationTest): + def setUpTensorRTParam(self): + self.enable_trt = True + self.trt_parameters = TensorRTSubgraphPassActivationTest.TensorRTParam( + 1 << 30, 32, 0, AnalysisConfig.Precision.Half, True, False) + + def append_act(self, x): + return fluid.layers.prelu(x, mode='all') + + +class TensorRTSubgraphPassPreluFp16DynamicTest( + TensorRTSubgraphPassActivationTest): + def setUpTensorRTParam(self): + self.enable_trt = True + self.trt_parameters = TensorRTSubgraphPassActivationTest.TensorRTParam( + 1 << 30, 32, 0, AnalysisConfig.Precision.Half, False, False) + self.dynamic_shape_params = TensorRTSubgraphPassActivationTest.DynamicShapeParam( + { + 'data': [1, 6, 8, 8] + }, {'data': [1, 6, 512, 512]}, {'data': [1, 6, 256, 256]}, False) + + def append_act(self, x): + return fluid.layers.prelu(x, mode='all') + + +class TensorRTSubgraphPassPreluFp16DynamicSerializeTest( + TensorRTSubgraphPassActivationTest): + def setUpTensorRTParam(self): + self.enable_trt = True + self.trt_parameters = TensorRTSubgraphPassActivationTest.TensorRTParam( + 1 << 30, 32, 0, AnalysisConfig.Precision.Half, True, False) + self.dynamic_shape_params = TensorRTSubgraphPassActivationTest.DynamicShapeParam( + { + 'data': [1, 6, 8, 8] + }, {'data': [1, 6, 512, 512]}, {'data': [1, 6, 256, 256]}, False) + + def append_act(self, x): + return fluid.layers.prelu(x, mode='all') + + class TensorRTSubgraphPassGeluTest(TensorRTSubgraphPassActivationTest): def append_act(self, x): return fluid.layers.gelu(x) From a340884afd70cc1852977cd22a27fbb8eb28c407 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Mon, 28 Jun 2021 12:25:04 +0000 Subject: [PATCH 40/46] fix pool dynamic plugin serialization and add ut --- .../tensorrt/plugin/pool_op_plugin.cu | 30 ++- .../tensorrt/plugin/pool_op_plugin.h | 13 +- .../ir/inference/test_trt_pool_op.py | 205 ++++++++++++++++++ .../ir/inference/test_trt_subgraph_pass.py | 107 --------- 4 files changed, 235 insertions(+), 120 deletions(-) create mode 100644 python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py diff --git a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu index 85fb495d64162..9d047942d4850 100644 --- a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu @@ -75,9 +75,35 @@ int PoolPlugin::enqueue(int batchSize, const void *const *inputs, // Dynamic Plugin below. #if IS_TRT_VERSION_GE(6000) -size_t PoolPluginDynamic::getSerializationSize() const { return 0; } +PoolPluginDynamic::PoolPluginDynamic(void const *serialData, + size_t serialLength) { + DeserializeValue(&serialData, &serialLength, &ceil_mode_); + const char *pool_type; + DeserializeValue(&serialData, &serialLength, &pool_type); + pool_type_ = std::string(pool_type); + DeserializeValue(&serialData, &serialLength, &adaptive_); + DeserializeValue(&serialData, &serialLength, &ksize_); + DeserializeValue(&serialData, &serialLength, &strides_); + DeserializeValue(&serialData, &serialLength, &paddings_); + DeserializeValue(&serialData, &serialLength, &is_global_); +} + +size_t PoolPluginDynamic::getSerializationSize() const { + return SerializedSize(ceil_mode_) + SerializedSize(pool_type_.c_str()) + + SerializedSize(adaptive_) + SerializedSize(ksize_) + + SerializedSize(strides_) + SerializedSize(paddings_) + + SerializedSize(is_global_); +} -void PoolPluginDynamic::serialize(void *buffer) const {} +void PoolPluginDynamic::serialize(void *buffer) const { + SerializeValue(&buffer, ceil_mode_); + SerializeValue(&buffer, pool_type_.c_str()); + SerializeValue(&buffer, adaptive_); + SerializeValue(&buffer, ksize_); + SerializeValue(&buffer, strides_); + SerializeValue(&buffer, paddings_); + SerializeValue(&buffer, is_global_); +} nvinfer1::DimsExprs PoolPluginDynamic::getOutputDimensions( int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs, diff --git a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h index 4c9f7c9e06dea..64fbfdc649c0e 100644 --- a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -174,17 +175,7 @@ class PoolPluginDynamic : public DynamicPluginTensorRT { paddings_(paddings), is_global_(is_global) {} - PoolPluginDynamic(void const* serialData, size_t serialLength) { - DeserializeValue(&serialData, &serialLength, &ceil_mode_); - const char* pool_type; - DeserializeValue(&serialData, &serialLength, &pool_type); - pool_type_ = std::string(pool_type); - DeserializeValue(&serialData, &serialLength, &adaptive_); - DeserializeValue(&serialData, &serialLength, &ksize_); - DeserializeValue(&serialData, &serialLength, &strides_); - DeserializeValue(&serialData, &serialLength, &paddings_); - DeserializeValue(&serialData, &serialLength, &is_global_); - } + PoolPluginDynamic(void const* serialData, size_t serialLength); ~PoolPluginDynamic() {} nvinfer1::IPluginV2DynamicExt* clone() const override { return new PoolPluginDynamic(ceil_mode_, pool_type_, adaptive_, ksize_, diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py new file mode 100644 index 0000000000000..2ddb713502b08 --- /dev/null +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py @@ -0,0 +1,205 @@ +# Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. +# Copyright (c) 2021 NVIDIA Corporation. All rights reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +import os +import shutil +import unittest +import numpy as np +from inference_pass_test import InferencePassTest +import paddle.fluid as fluid +import paddle.fluid.core as core +from paddle.fluid.core import PassVersionChecker +from paddle.fluid.core import AnalysisConfig + + +class TensorRTPoolTest(InferencePassTest): + def setUp(self): + self.bs = 1 + self.channel = 4 + self.height = 16 + self.width = 16 + self.pool_size = 2 + self.pool_type = 'max' + self.pool_stride = 1 + self.pool_padding = 0 + self.global_pooling = False + self.ceil_mode = False + self.exclusive = False + self.enable_trt = True + self.serialize = False + self.precision = AnalysisConfig.Precision.Float32 + self.feeds = { + 'data': + np.random.random([self.bs, self.channel, self.height, + self.width]).astype('float32'), + } + + def set_extra_config(self): + pass + + def build_network(self): + self.set_extra_config() + self.trt_parameters = TensorRTPoolTest.TensorRTParam( + 1 << 30, self.bs, 0, self.precision, self.serialize, False) + + with fluid.program_guard(self.main_program, self.startup_program): + data = fluid.data( + name='data', + shape=[-1, self.channel, self.height, self.width], + dtype='float32') + pool_out = fluid.layers.pool2d( + input=data, + pool_size=self.pool_size, + pool_type=self.pool_type, + pool_stride=self.pool_stride, + pool_padding=self.pool_padding, + global_pooling=self.global_pooling, + ceil_mode=self.ceil_mode, + exclusive=self.exclusive) + out = fluid.layers.batch_norm(pool_out, is_test=True) + self.fetch_list = [out] + + def check_output(self): + if os.path.exists(self.path + "_opt_cache"): + shutil.rmtree(self.path + "_opt_cache") + if core.is_compiled_with_cuda(): + use_gpu = True + self.check_output_with_option(use_gpu) + self.assertTrue( + PassVersionChecker.IsCompatible('tensorrt_subgraph_pass')) + + def set_dynamic(self): + self.dynamic_shape_params = InferencePassTest.DynamicShapeParam( + { + 'data': + [self.bs, self.channel, self.height // 2, self.width // 2] + }, {'data': [self.bs, self.channel, self.height, self.width]}, + {'data': [self.bs, self.channel, self.height, self.width]}, False) + + def set_fp16(self): + self.precision = AnalysisConfig.Precision.Half + + def set_serialize(self): + self.serialize = True + + def run_test(self): + self.build_network() + self.check_output() + + def test_fp32(self): + self.run_test() + + def test_fp16(self): + self.set_fp16() + self.run_test() + + def test_fp32_serialize(self): + self.set_serialize() + self.run_test() + + def test_fp16_serialize(self): + self.set_fp16() + self.set_serialize() + self.run_test() + + def test_fp32_dynamic(self): + self.set_dynamic() + self.run_test() + + def test_fp16_dynamic(self): + self.set_fp16() + self.set_dynamic() + self.run_test() + + def test_fp32_dynamic_serialize(self): + self.set_dynamic() + self.set_serialize() + self.run_test() + + def test_fp16_dynamic_serialize(self): + self.set_fp16() + self.set_dynamic() + self.set_serialize() + self.run_test() + + +class TensorRTAvgPoolTest(TensorRTPoolTest): + def set_extra_config(self): + self.pool_size = 2 + self.pool_type = 'avg' + self.pool_stride = 1 + self.pool_padding = 0 + self.global_pooling = False + self.ceil_mode = False + self.exclusive = False + + +class TensorRTGlobalPoolTest(TensorRTPoolTest): + def set_extra_config(self): + self.pool_size = 2 + self.pool_type = 'max' + self.pool_stride = 1 + self.pool_padding = 0 + self.global_pooling = True + self.ceil_mode = False + self.exclusive = False + + +class TensorRTCeilPoolTest(TensorRTPoolTest): + def set_extra_config(self): + self.pool_size = 2 + self.pool_type = 'max' + self.pool_stride = 1 + self.pool_padding = 0 + self.global_pooling = False + self.ceil_mode = True + self.exclusive = False + + +class TensorRTExclusivePoolTest(TensorRTPoolTest): + def set_extra_config(self): + self.pool_size = 2 + self.pool_type = 'max' + self.pool_stride = 1 + self.pool_padding = 0 + self.global_pooling = False + self.ceil_mode = False + self.exclusive = True + + +class TensorRTSamePaddingPoolTest(InferencePassTest): + def set_extra_config(self): + self.pool_size = 2 + self.pool_type = 'max' + self.pool_stride = 1 + self.pool_padding = 'SAME' + self.global_pooling = False + self.ceil_mode = False + self.exclusive = False + + +class TensorRTValidPaddingPoolTest(InferencePassTest): + def set_extra_config(self): + self.pool_size = 2 + self.pool_type = 'max' + self.pool_stride = 1 + self.pool_padding = 'VALID' + self.global_pooling = False + self.ceil_mode = False + self.exclusive = False + + +if __name__ == "__main__": + unittest.main() diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py index d85f705c88135..23a3d19140179 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py @@ -47,113 +47,6 @@ def test_check_output(self): PassVersionChecker.IsCompatible('tensorrt_subgraph_pass')) -class TensorRTSubgraphPassPoolTest(InferencePassTest): - def setUp(self): - self.set_params() - with fluid.program_guard(self.main_program, self.startup_program): - data = fluid.data( - name="data", shape=[-1, 6, 64, 64], dtype="float32") - pool_out = fluid.layers.pool2d( - input=data, - pool_size=self.pool_size, - pool_type=self.pool_type, - pool_stride=self.pool_stride, - pool_padding=self.pool_padding, - global_pooling=self.global_pooling, - ceil_mode=self.ceil_mode, - exclusive=self.exclusive) - out = fluid.layers.batch_norm(pool_out, is_test=True) - self.feeds = { - "data": np.random.random([1, 6, 64, 64]).astype("float32"), - } - self.enable_trt = True - self.trt_parameters = TensorRTSubgraphPassPoolTest.TensorRTParam( - 1 << 30, 32, 0, AnalysisConfig.Precision.Float32, False, False) - self.fetch_list = [out] - - def set_params(self): - self.pool_size = 2 - self.pool_type = 'max' - self.pool_stride = 1 - self.pool_padding = 0 - self.global_pooling = False - self.ceil_mode = False - self.exclusive = False - - def test_check_output(self): - if core.is_compiled_with_cuda(): - use_gpu = True - self.check_output_with_option(use_gpu) - self.assertTrue( - PassVersionChecker.IsCompatible('tensorrt_subgraph_pass')) - - -class TensorRTSubgraphPassAvgPoolTest(TensorRTSubgraphPassPoolTest): - def set_params(self): - self.pool_size = 2 - self.pool_type = 'avg' - self.pool_stride = 1 - self.pool_padding = 0 - self.global_pooling = False - self.ceil_mode = False - self.exclusive = False - - -class TensorRTSubgraphPassGlobalPoolTest(TensorRTSubgraphPassPoolTest): - def set_params(self): - self.pool_size = 2 - self.pool_type = 'max' - self.pool_stride = 1 - self.pool_padding = 0 - self.global_pooling = True - self.ceil_mode = False - self.exclusive = False - - -class TensorRTSubgraphPassCeilPoolTest(TensorRTSubgraphPassPoolTest): - def set_params(self): - self.pool_size = 2 - self.pool_type = 'max' - self.pool_stride = 1 - self.pool_padding = 0 - self.global_pooling = False - self.ceil_mode = True - self.exclusive = False - - -class TensorRTSubgraphPassExclusivePoolTest(TensorRTSubgraphPassPoolTest): - def set_params(self): - self.pool_size = 2 - self.pool_type = 'max' - self.pool_stride = 1 - self.pool_padding = 0 - self.global_pooling = False - self.ceil_mode = False - self.exclusive = True - - -class TensorRTSubgraphPassSamePaddingPoolTest(InferencePassTest): - def set_params(self): - self.pool_size = 2 - self.pool_type = 'max' - self.pool_stride = 1 - self.pool_padding = 'SAME' - self.global_pooling = False - self.ceil_mode = False - self.exclusive = False - - -class TensorRTSubgraphPassValidPaddingPoolTest(InferencePassTest): - def set_params(self): - self.pool_size = 2 - self.pool_type = 'max' - self.pool_stride = 1 - self.pool_padding = 'VALID' - self.global_pooling = False - self.ceil_mode = False - self.exclusive = False - - class TensorRTSubgraphPassConcatTest(InferencePassTest): def setUp(self): with fluid.program_guard(self.main_program, self.startup_program): From 6097f424b4579c951003f262a84f71e01b5bf805 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Mon, 28 Jun 2021 13:04:19 +0000 Subject: [PATCH 41/46] refine pool ut with subtest --- .../ir/inference/test_trt_pool_op.py | 70 ++++++------------- 1 file changed, 22 insertions(+), 48 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py index 2ddb713502b08..648687043f554 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py @@ -16,6 +16,7 @@ import os import shutil import unittest +import itertools import numpy as np from inference_pass_test import InferencePassTest import paddle.fluid as fluid @@ -80,59 +81,32 @@ def check_output(self): self.assertTrue( PassVersionChecker.IsCompatible('tensorrt_subgraph_pass')) - def set_dynamic(self): - self.dynamic_shape_params = InferencePassTest.DynamicShapeParam( + def run_test(self): + self.build_network() + self.check_output() + + def test(self): + precision_options = [ + AnalysisConfig.Precision.Float32, AnalysisConfig.Precision.Half + ] + serialize_options = [False, True] + dynamic_shape_profile = InferencePassTest.DynamicShapeParam( { 'data': [self.bs, self.channel, self.height // 2, self.width // 2] }, {'data': [self.bs, self.channel, self.height, self.width]}, {'data': [self.bs, self.channel, self.height, self.width]}, False) - - def set_fp16(self): - self.precision = AnalysisConfig.Precision.Half - - def set_serialize(self): - self.serialize = True - - def run_test(self): - self.build_network() - self.check_output() - - def test_fp32(self): - self.run_test() - - def test_fp16(self): - self.set_fp16() - self.run_test() - - def test_fp32_serialize(self): - self.set_serialize() - self.run_test() - - def test_fp16_serialize(self): - self.set_fp16() - self.set_serialize() - self.run_test() - - def test_fp32_dynamic(self): - self.set_dynamic() - self.run_test() - - def test_fp16_dynamic(self): - self.set_fp16() - self.set_dynamic() - self.run_test() - - def test_fp32_dynamic_serialize(self): - self.set_dynamic() - self.set_serialize() - self.run_test() - - def test_fp16_dynamic_serialize(self): - self.set_fp16() - self.set_dynamic() - self.set_serialize() - self.run_test() + dynamic_shape_options = [None, dynamic_shape_profile] + + for precision, serialize, dynamic_shape in itertools.product( + precision_options, serialize_options, dynamic_shape_options): + is_dynamic = True if dynamic_shape_options is not None else False + with self.subTest('Precision: {}, Serialize: {}, Dynamic: {}'. + format(precision, serialize, is_dynamic)): + self.precision = precision + self.serialize = serialize + self.dynamic_shape = dynamic_shape + self.run_test() class TensorRTAvgPoolTest(TensorRTPoolTest): From eafd8c758ab74b538b341d23808f9e1974220cee Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Mon, 28 Jun 2021 13:04:48 +0000 Subject: [PATCH 42/46] add env for avoiding oom --- python/paddle/fluid/tests/unittests/ir/inference/CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/python/paddle/fluid/tests/unittests/ir/inference/CMakeLists.txt b/python/paddle/fluid/tests/unittests/ir/inference/CMakeLists.txt index 0f068045e0c09..905cc18466d24 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/CMakeLists.txt +++ b/python/paddle/fluid/tests/unittests/ir/inference/CMakeLists.txt @@ -35,4 +35,5 @@ set_tests_properties(test_trt_activation_pass PROPERTIES TIMEOUT 120) set_tests_properties(test_trt_conv_pass PROPERTIES TIMEOUT 120) #set_tests_properties(test_trt_multiclass_nms_op PROPERTIES TIMEOUT 200) set_tests_properties(test_trt_dynamic_shape PROPERTIES TIMEOUT 120) +set_tests_properties(test_trt_pool_op PROPERTIES ENVIRONMENT FLAGS_fraction_of_gpu_memory_to_use=0.1) endif() From d9a6505f835774789f9cb1124d0231ce77d94977 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Mon, 28 Jun 2021 15:26:07 +0000 Subject: [PATCH 43/46] reduce test input size & increase pool op ut to 45s --- .../fluid/tests/unittests/ir/inference/CMakeLists.txt | 2 +- .../fluid/tests/unittests/ir/inference/test_trt_pool_op.py | 6 +++--- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/ir/inference/CMakeLists.txt b/python/paddle/fluid/tests/unittests/ir/inference/CMakeLists.txt index 905cc18466d24..792a976aeb025 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/CMakeLists.txt +++ b/python/paddle/fluid/tests/unittests/ir/inference/CMakeLists.txt @@ -35,5 +35,5 @@ set_tests_properties(test_trt_activation_pass PROPERTIES TIMEOUT 120) set_tests_properties(test_trt_conv_pass PROPERTIES TIMEOUT 120) #set_tests_properties(test_trt_multiclass_nms_op PROPERTIES TIMEOUT 200) set_tests_properties(test_trt_dynamic_shape PROPERTIES TIMEOUT 120) -set_tests_properties(test_trt_pool_op PROPERTIES ENVIRONMENT FLAGS_fraction_of_gpu_memory_to_use=0.1) +set_tests_properties(test_trt_pool_op PROPERTIES ENVIRONMENT FLAGS_fraction_of_gpu_memory_to_use=0.1 TIMEOUT 45) endif() diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py index 648687043f554..5204715bf51b9 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py @@ -28,9 +28,9 @@ class TensorRTPoolTest(InferencePassTest): def setUp(self): self.bs = 1 - self.channel = 4 - self.height = 16 - self.width = 16 + self.channel = 3 + self.height = 8 + self.width = 8 self.pool_size = 2 self.pool_type = 'max' self.pool_stride = 1 From 80cef1e50a381c83e86d36832a2000590b6b3181 Mon Sep 17 00:00:00 2001 From: zlsh80826 Date: Tue, 29 Jun 2021 02:42:32 +0000 Subject: [PATCH 44/46] add the contributor --- paddle/fluid/inference/tensorrt/convert/elementwise_op.cc | 1 + paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc | 1 + paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc | 1 + paddle/fluid/inference/tensorrt/engine.cc | 1 + paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h | 1 + paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h | 1 + paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h | 1 + paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h | 1 + paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h | 1 + paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h | 1 + paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h | 1 + paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc | 1 + paddle/fluid/inference/tensorrt/plugin/trt_plugin.h | 1 + .../tests/unittests/ir/inference/test_trt_activation_pass.py | 1 + .../tests/unittests/ir/inference/test_trt_elementwise_op.py | 1 + .../tests/unittests/ir/inference/test_trt_instance_norm_op.py | 1 + .../fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py | 1 + 17 files changed, 17 insertions(+) diff --git a/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc b/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc index 2f802ea8d181e..5a4f8d37de649 100644 --- a/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc +++ b/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc @@ -1,4 +1,5 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. + * Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc b/paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc index b7097fc05680d..b6e86bbff980b 100644 --- a/paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc +++ b/paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc @@ -1,4 +1,5 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. + * Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc b/paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc index 976fe9502acd6..42114c855f2f1 100644 --- a/paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc +++ b/paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc @@ -1,4 +1,5 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. + * Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/engine.cc b/paddle/fluid/inference/tensorrt/engine.cc index e77e12713ca20..3b01d24c24a75 100644 --- a/paddle/fluid/inference/tensorrt/engine.cc +++ b/paddle/fluid/inference/tensorrt/engine.cc @@ -1,4 +1,5 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. + * Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h index 5dd3142c75839..d2dbaa07cdb52 100644 --- a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h @@ -1,4 +1,5 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. + * Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h index 6fdd9791a61bd..d626431fb183b 100644 --- a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h @@ -1,4 +1,5 @@ // Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h index 42c47959988a5..b93bdc3fe3748 100644 --- a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h @@ -1,4 +1,5 @@ // Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h index f9dab09beebd3..f38a2896fce75 100644 --- a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h index caa3c21db63fa..2e13dfc3f9f9f 100644 --- a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h index b656918f8fbab..5455721101d78 100644 --- a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h index 8940fdce3b0b5..2c1bd656ace1a 100644 --- a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc index 5be0ed4a13b23..48c804f401903 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h index 599294392799d..14ce49d3d2363 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h @@ -1,4 +1,5 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. +// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_activation_pass.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_activation_pass.py index 8e196f5081f73..c2aa51bb64753 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_activation_pass.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_activation_pass.py @@ -1,4 +1,5 @@ # Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. +# Copyright (c) 2021 NVIDIA Corporation. All rights reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py index f84202df5fb93..f4479f7140651 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py @@ -1,4 +1,5 @@ # Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. +# Copyright (c) 2021 NVIDIA Corporation. All rights reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py index d283465dcba09..a17a9b5bb53c0 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py @@ -1,4 +1,5 @@ # Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. +# Copyright (c) 2021 NVIDIA Corporation. All rights reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py index 23a3d19140179..21ab35d23fd69 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py @@ -1,4 +1,5 @@ # Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. +# Copyright (c) 2021 NVIDIA Corporation. All rights reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. From ca6b8b3899366250ea4f42fb7ea9d109b50ab0f4 Mon Sep 17 00:00:00 2001 From: zlhs80826 Date: Thu, 8 Jul 2021 22:59:56 +0800 Subject: [PATCH 45/46] remove copyright (will add in contributor) --- paddle/fluid/inference/tensorrt/convert/elementwise_op.cc | 1 - paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc | 1 - paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc | 1 - paddle/fluid/inference/tensorrt/engine.cc | 1 - paddle/fluid/inference/tensorrt/engine.h | 1 - .../inference/tensorrt/plugin/anchor_generator_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h | 1 - .../inference/tensorrt/plugin/emb_eltwise_layernorm_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/gather_nd_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h | 1 - paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h | 1 - .../fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h | 1 - paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h | 1 - paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h | 1 - paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h | 1 - paddle/fluid/inference/tensorrt/plugin/qkv_to_context_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/roi_align_op_plugin.cu | 1 - .../fluid/inference/tensorrt/plugin/skip_layernorm_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h | 1 - paddle/fluid/inference/tensorrt/plugin/special_slice_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/stack_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu | 1 - paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h | 1 - paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc | 1 - paddle/fluid/inference/tensorrt/plugin/trt_plugin.h | 1 - paddle/fluid/inference/tensorrt/plugin/yolo_box_op_plugin.cu | 1 - 35 files changed, 35 deletions(-) diff --git a/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc b/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc index 5a4f8d37de649..2f802ea8d181e 100644 --- a/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc +++ b/paddle/fluid/inference/tensorrt/convert/elementwise_op.cc @@ -1,5 +1,4 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. - * Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc b/paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc index b6e86bbff980b..b7097fc05680d 100644 --- a/paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc +++ b/paddle/fluid/inference/tensorrt/convert/instance_norm_op.cc @@ -1,5 +1,4 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. - * Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc b/paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc index 42114c855f2f1..976fe9502acd6 100644 --- a/paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc +++ b/paddle/fluid/inference/tensorrt/convert/shuffle_channel_op.cc @@ -1,5 +1,4 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. - * Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/engine.cc b/paddle/fluid/inference/tensorrt/engine.cc index 3b01d24c24a75..e77e12713ca20 100644 --- a/paddle/fluid/inference/tensorrt/engine.cc +++ b/paddle/fluid/inference/tensorrt/engine.cc @@ -1,5 +1,4 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. - * Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/engine.h b/paddle/fluid/inference/tensorrt/engine.h index 72d68f1cd7f1a..38c453bde6d2d 100644 --- a/paddle/fluid/inference/tensorrt/engine.h +++ b/paddle/fluid/inference/tensorrt/engine.h @@ -1,5 +1,4 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. - Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/anchor_generator_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/anchor_generator_op_plugin.cu index cadfa1185c5be..29fab72733337 100644 --- a/paddle/fluid/inference/tensorrt/plugin/anchor_generator_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/anchor_generator_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu index 5bbab49eaf558..3338aae370e51 100644 --- a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu @@ -1,5 +1,4 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. - Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h index d2dbaa07cdb52..5dd3142c75839 100644 --- a/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h @@ -1,5 +1,4 @@ /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. - * Copyright (c) 2021 NVIDIA Corporation. All rights reserved. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/emb_eltwise_layernorm_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/emb_eltwise_layernorm_plugin.cu index 9105619660189..79fc3d66bbe4d 100644 --- a/paddle/fluid/inference/tensorrt/plugin/emb_eltwise_layernorm_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/emb_eltwise_layernorm_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/gather_nd_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/gather_nd_op_plugin.cu index c9d30f7ccf547..933ca333cdbb9 100644 --- a/paddle/fluid/inference/tensorrt/plugin/gather_nd_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/gather_nd_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2021 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu index 2275c742d6d28..43557c341ef42 100644 --- a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h index d626431fb183b..6fdd9791a61bd 100644 --- a/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h @@ -1,5 +1,4 @@ // Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu index 3e37f04370d76..dab7ddac1957a 100644 --- a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h index b93bdc3fe3748..42c47959988a5 100644 --- a/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h @@ -1,5 +1,4 @@ // Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu index 4d7720d69d1a7..13aa6df643e82 100644 --- a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h index f38a2896fce75..f9dab09beebd3 100644 --- a/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu index a2fbfe1f7d74f..2688380726f78 100644 --- a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h index 2e13dfc3f9f9f..caa3c21db63fa 100644 --- a/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu index 9d047942d4850..7e1d18227e232 100644 --- a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h index 64fbfdc649c0e..7c12796805c5d 100644 --- a/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu index a707f7c030ce5..1882084a8f516 100644 --- a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h index 3bc126f75224b..e3f05bdbe85a1 100644 --- a/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/qkv_to_context_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/qkv_to_context_plugin.cu index c9ae888a170a0..0d9e5417263f3 100644 --- a/paddle/fluid/inference/tensorrt/plugin/qkv_to_context_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/qkv_to_context_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/roi_align_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/roi_align_op_plugin.cu index 39feded7e1ae6..c08c0527072bf 100644 --- a/paddle/fluid/inference/tensorrt/plugin/roi_align_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/roi_align_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/skip_layernorm_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/skip_layernorm_op_plugin.cu index 6a8922f6b6b97..346b4c680830e 100644 --- a/paddle/fluid/inference/tensorrt/plugin/skip_layernorm_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/skip_layernorm_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu index 862ddc009eb4b..70ff0e7cb069d 100644 --- a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h index 5455721101d78..b656918f8fbab 100644 --- a/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/special_slice_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/special_slice_plugin.cu index b7c76335ccb8e..3bef9672e5058 100644 --- a/paddle/fluid/inference/tensorrt/plugin/special_slice_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/special_slice_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu index 85ae595e7588f..37afff9105d80 100644 --- a/paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/stack_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/stack_op_plugin.cu index c796990d9ef1d..21e80339b5006 100644 --- a/paddle/fluid/inference/tensorrt/plugin/stack_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/stack_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu index 70b2e48c23d21..da9d21acd5d63 100644 --- a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h index 2c1bd656ace1a..8940fdce3b0b5 100644 --- a/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc index 48c804f401903..5be0ed4a13b23 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h index 14ce49d3d2363..599294392799d 100644 --- a/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h +++ b/paddle/fluid/inference/tensorrt/plugin/trt_plugin.h @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. diff --git a/paddle/fluid/inference/tensorrt/plugin/yolo_box_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/yolo_box_op_plugin.cu index de635a7f08530..490c42b621b25 100644 --- a/paddle/fluid/inference/tensorrt/plugin/yolo_box_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/yolo_box_op_plugin.cu @@ -1,5 +1,4 @@ // Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. -// Copyright (c) 2021 NVIDIA Corporation. All rights reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. From df49070c94ca88dbdecf13f69dfc40909a39c7e2 Mon Sep 17 00:00:00 2001 From: zlhs80826 Date: Fri, 9 Jul 2021 14:14:51 +0800 Subject: [PATCH 46/46] remove copyright (will add in contributor) --- .../tests/unittests/ir/inference/test_trt_activation_pass.py | 1 - .../tests/unittests/ir/inference/test_trt_elementwise_op.py | 1 - .../tests/unittests/ir/inference/test_trt_instance_norm_op.py | 1 - .../fluid/tests/unittests/ir/inference/test_trt_pool_op.py | 1 - .../fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py | 1 - 5 files changed, 5 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_activation_pass.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_activation_pass.py index c2aa51bb64753..8e196f5081f73 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_activation_pass.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_activation_pass.py @@ -1,5 +1,4 @@ # Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. -# Copyright (c) 2021 NVIDIA Corporation. All rights reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py index f4479f7140651..f84202df5fb93 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_elementwise_op.py @@ -1,5 +1,4 @@ # Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. -# Copyright (c) 2021 NVIDIA Corporation. All rights reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py index a17a9b5bb53c0..d283465dcba09 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_instance_norm_op.py @@ -1,5 +1,4 @@ # Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. -# Copyright (c) 2021 NVIDIA Corporation. All rights reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py index 5204715bf51b9..3d317446f00f3 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_pool_op.py @@ -1,5 +1,4 @@ # Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. -# Copyright (c) 2021 NVIDIA Corporation. All rights reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. diff --git a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py index 21ab35d23fd69..23a3d19140179 100644 --- a/python/paddle/fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py +++ b/python/paddle/fluid/tests/unittests/ir/inference/test_trt_subgraph_pass.py @@ -1,5 +1,4 @@ # Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved. -# Copyright (c) 2021 NVIDIA Corporation. All rights reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License.