Skip to content

P
Paddle

BaiXuePrincess / Paddle
与 Fork 源项目一致

Fork自 PaddlePaddle / Paddle

通知 1
Star 1
Fork 0
P
Paddle
未验证 提交 b5aab4f0 编写于 作者: Z zlsh80826 提交者： GitHub
浏览文件
操作

[Paddle-TRT] Add noexcept on methods inherited from TensorRT (#34157) 

* add trt noexcept definition

* add trt noexcept on trt plugin

* add trt noexcept on trt int8 calibrator

* remove noexcept on base serialize

* add trt noexcept on split plugin

* add trt noexcept on elementwise plugin

* add trt noexcept on prelu plugin

* add trt noexcept on pool plugin

* add trt noexcept on swish plugin

* add trt noexcept on gelu plugin

* add trt noexcept on layer norm plugin

* add trt noexcept on instance norm plugin

* add trt noexcept on emb eltwise layernorm plugin

* add trt noexcept on qkv2context plugin

* add trt noexcept on skip layernorm plugin

* add trt noexcept on slice plugin

* add trt noexcept on hard swish plugin

* add trt noexcept on stack plugin

* add trt noexcept on special slice plugin

* add trt noexcept on anchor generator plugin

* add trt noexcept on yolobox plugin

* add trt noexcept on roi align plugin

* add trt noexcept on gather nd plugin
上级 c0133e01
隐藏空白更改
内联 并排
Showing with 1143 addition and 942 deletion
paddle/fluid/inference/tensorrt/helper.h
浏览文件 @ b5aab4f0
...@@ -39,6 +39,12 @@ namespace tensorrt { ...@@ -39,6 +39,12 @@ namespace tensorrt {
NV_TENSORRT_MAJOR * 1000 + NV_TENSORRT_MINOR * 100 + \ NV_TENSORRT_MAJOR * 1000 + NV_TENSORRT_MINOR * 100 + \
NV_TENSORRT_PATCH * 10 + NV_TENSORRT_BUILD NV_TENSORRT_PATCH * 10 + NV_TENSORRT_BUILD
#if IS_TRT_VERSION_GE(8000)
#define TRT_NOEXCEPT noexcept
#else
#define TRT_NOEXCEPT
#endif
namespace dy = paddle::platform::dynload; namespace dy = paddle::platform::dynload;
// TensorRT data type to size // TensorRT data type to size
...@@ -72,7 +78,8 @@ static int GetInferLibVersion() { ...@@ -72,7 +78,8 @@ static int GetInferLibVersion() {
// A logger for create TensorRT infer builder. // A logger for create TensorRT infer builder.
class NaiveLogger : public nvinfer1::ILogger { class NaiveLogger : public nvinfer1::ILogger {
public: public:
void log(nvinfer1::ILogger::Severity severity, const char* msg) override { void log(nvinfer1::ILogger::Severity severity,
const char* msg) TRT_NOEXCEPT override {
switch (severity) { switch (severity) {
case Severity::kVERBOSE: case Severity::kVERBOSE:
VLOG(3) << msg; VLOG(3) << msg;
...@@ -105,7 +112,7 @@ class NaiveProfiler : public nvinfer1::IProfiler { ...@@ -105,7 +112,7 @@ class NaiveProfiler : public nvinfer1::IProfiler {
typedef std::pair<std::string, float> Record; typedef std::pair<std::string, float> Record;
std::vector<Record> mProfile; std::vector<Record> mProfile;
virtual void reportLayerTime(const char* layerName, float ms) { virtual void reportLayerTime(const char* layerName, float ms) TRT_NOEXCEPT {
auto record = auto record =
std::find_if(mProfile.begin(), mProfile.end(), std::find_if(mProfile.begin(), mProfile.end(),
[&](const Record& r) { return r.first == layerName; }); [&](const Record& r) { return r.first == layerName; });
......
paddle/fluid/inference/tensorrt/plugin/anchor_generator_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -110,16 +110,18 @@ AnchorGeneratorPlugin::AnchorGeneratorPlugin(const void* data, size_t length) { ...@@ -110,16 +110,18 @@ AnchorGeneratorPlugin::AnchorGeneratorPlugin(const void* data, size_t length) {
PrepareParamsOnDevice(); PrepareParamsOnDevice();
} }
const char* AnchorGeneratorPlugin::getPluginType() const { const char* AnchorGeneratorPlugin::getPluginType() const TRT_NOEXCEPT {
return "anchor_generator_plugin"; return "anchor_generator_plugin";
} }
const char* AnchorGeneratorPlugin::getPluginVersion() const { return "1"; } const char* AnchorGeneratorPlugin::getPluginVersion() const TRT_NOEXCEPT {
return "1";
}
int AnchorGeneratorPlugin::getNbOutputs() const { return 2; } int AnchorGeneratorPlugin::getNbOutputs() const TRT_NOEXCEPT { return 2; }
nvinfer1::Dims AnchorGeneratorPlugin::getOutputDimensions( nvinfer1::Dims AnchorGeneratorPlugin::getOutputDimensions(
int index, const nvinfer1::Dims* inputs, int nb_input_dims) { int index, const nvinfer1::Dims* inputs, int nb_input_dims) TRT_NOEXCEPT {
nvinfer1::Dims dims{}; nvinfer1::Dims dims{};
dims.nbDims = 4; dims.nbDims = 4;
dims.d[0] = height_; dims.d[0] = height_;
...@@ -130,20 +132,21 @@ nvinfer1::Dims AnchorGeneratorPlugin::getOutputDimensions( ...@@ -130,20 +132,21 @@ nvinfer1::Dims AnchorGeneratorPlugin::getOutputDimensions(
} }
bool AnchorGeneratorPlugin::supportsFormat( bool AnchorGeneratorPlugin::supportsFormat(
nvinfer1::DataType type, nvinfer1::TensorFormat format) const { nvinfer1::DataType type, nvinfer1::TensorFormat format) const TRT_NOEXCEPT {
// static shape plugin can't support different type between input/out // static shape plugin can't support different type between input/out
// it may cause addition overhead in half mode // it may cause addition overhead in half mode
return (type == data_type_ && format == nvinfer1::TensorFormat::kLINEAR); return (type == data_type_ && format == nvinfer1::TensorFormat::kLINEAR);
} }
size_t AnchorGeneratorPlugin::getWorkspaceSize(int max_batch_size) const { size_t AnchorGeneratorPlugin::getWorkspaceSize(int max_batch_size) const
TRT_NOEXCEPT {
return 0; return 0;
} }
template <typename T> template <typename T>
int AnchorGeneratorPlugin::enqueue_impl(int batch_size, int AnchorGeneratorPlugin::enqueue_impl(int batch_size,
const void* const* inputs, const void* const* inputs,
void** outputs, void* workspace, void* const* outputs, void* workspace,
cudaStream_t stream) { cudaStream_t stream) {
const int block = 512; const int block = 512;
const int gen_anchor_grid = (box_num_ + block - 1) / block; const int gen_anchor_grid = (box_num_ + block - 1) / block;
...@@ -169,15 +172,15 @@ int AnchorGeneratorPlugin::enqueue(int batch_size, const void* const* inputs, ...@@ -169,15 +172,15 @@ int AnchorGeneratorPlugin::enqueue(int batch_size, const void* const* inputs,
#else #else
void* const* outputs, void* workspace, void* const* outputs, void* workspace,
#endif #endif
cudaStream_t stream) { cudaStream_t stream) TRT_NOEXCEPT {
return enqueue_impl<float>(batch_size, inputs, outputs, workspace, stream); return enqueue_impl<float>(batch_size, inputs, outputs, workspace, stream);
} }
int AnchorGeneratorPlugin::initialize() { return 0; } int AnchorGeneratorPlugin::initialize() TRT_NOEXCEPT { return 0; }
void AnchorGeneratorPlugin::terminate() {} void AnchorGeneratorPlugin::terminate() TRT_NOEXCEPT {}
size_t AnchorGeneratorPlugin::getSerializationSize() const { size_t AnchorGeneratorPlugin::getSerializationSize() const TRT_NOEXCEPT {
size_t serialize_size = 0; size_t serialize_size = 0;
serialize_size += SerializedSize(data_type_); serialize_size += SerializedSize(data_type_);
serialize_size += SerializedSize(anchor_sizes_); serialize_size += SerializedSize(anchor_sizes_);
...@@ -192,7 +195,7 @@ size_t AnchorGeneratorPlugin::getSerializationSize() const { ...@@ -192,7 +195,7 @@ size_t AnchorGeneratorPlugin::getSerializationSize() const {
return serialize_size; return serialize_size;
} }
void AnchorGeneratorPlugin::serialize(void* buffer) const { void AnchorGeneratorPlugin::serialize(void* buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, data_type_); SerializeValue(&buffer, data_type_);
SerializeValue(&buffer, anchor_sizes_); SerializeValue(&buffer, anchor_sizes_);
SerializeValue(&buffer, aspect_ratios_); SerializeValue(&buffer, aspect_ratios_);
...@@ -205,28 +208,31 @@ void AnchorGeneratorPlugin::serialize(void* buffer) const { ...@@ -205,28 +208,31 @@ void AnchorGeneratorPlugin::serialize(void* buffer) const {
SerializeValue(&buffer, box_num_); SerializeValue(&buffer, box_num_);
} }
void AnchorGeneratorPlugin::destroy() {} void AnchorGeneratorPlugin::destroy() TRT_NOEXCEPT {}
void AnchorGeneratorPlugin::setPluginNamespace(const char* lib_namespace) { void AnchorGeneratorPlugin::setPluginNamespace(const char* lib_namespace)
TRT_NOEXCEPT {
namespace_ = std::string(lib_namespace); namespace_ = std::string(lib_namespace);
} }
const char* AnchorGeneratorPlugin::getPluginNamespace() const { const char* AnchorGeneratorPlugin::getPluginNamespace() const TRT_NOEXCEPT {
return namespace_.c_str(); return namespace_.c_str();
} }
nvinfer1::DataType AnchorGeneratorPlugin::getOutputDataType( nvinfer1::DataType AnchorGeneratorPlugin::getOutputDataType(
int index, const nvinfer1::DataType* input_type, int nb_inputs) const { int index, const nvinfer1::DataType* input_type,
int nb_inputs) const TRT_NOEXCEPT {
return input_type[0]; return input_type[0];
} }
bool AnchorGeneratorPlugin::isOutputBroadcastAcrossBatch( bool AnchorGeneratorPlugin::isOutputBroadcastAcrossBatch(
int output_index, const bool* input_is_broadcast, int nb_inputs) const { int output_index, const bool* input_is_broadcast,
int nb_inputs) const TRT_NOEXCEPT {
return true; return true;
} }
bool AnchorGeneratorPlugin::canBroadcastInputAcrossBatch( bool AnchorGeneratorPlugin::canBroadcastInputAcrossBatch(int input_index) const
int input_index) const { TRT_NOEXCEPT {
return false; return false;
} }
...@@ -236,9 +242,9 @@ void AnchorGeneratorPlugin::configurePlugin( ...@@ -236,9 +242,9 @@ void AnchorGeneratorPlugin::configurePlugin(
const nvinfer1::DataType* input_types, const nvinfer1::DataType* input_types,
const nvinfer1::DataType* output_types, const bool* input_is_broadcast, const nvinfer1::DataType* output_types, const bool* input_is_broadcast,
const bool* output_is_broadcast, nvinfer1::PluginFormat float_format, const bool* output_is_broadcast, nvinfer1::PluginFormat float_format,
int max_batct_size) {} int max_batct_size) TRT_NOEXCEPT {}
nvinfer1::IPluginV2Ext* AnchorGeneratorPlugin::clone() const { nvinfer1::IPluginV2Ext* AnchorGeneratorPlugin::clone() const TRT_NOEXCEPT {
auto plugin = new AnchorGeneratorPlugin( auto plugin = new AnchorGeneratorPlugin(
data_type_, anchor_sizes_, aspect_ratios_, stride_, variances_, offset_, data_type_, anchor_sizes_, aspect_ratios_, stride_, variances_, offset_,
height_, width_, num_anchors_, box_num_); height_, width_, num_anchors_, box_num_);
...@@ -246,30 +252,32 @@ nvinfer1::IPluginV2Ext* AnchorGeneratorPlugin::clone() const { ...@@ -246,30 +252,32 @@ nvinfer1::IPluginV2Ext* AnchorGeneratorPlugin::clone() const {
return plugin; return plugin;
} }
void AnchorGeneratorPluginCreator::setPluginNamespace( void AnchorGeneratorPluginCreator::setPluginNamespace(const char* lib_namespace)
const char* lib_namespace) { TRT_NOEXCEPT {
namespace_ = std::string(lib_namespace); namespace_ = std::string(lib_namespace);
} }
const char* AnchorGeneratorPluginCreator::getPluginNamespace() const { const char* AnchorGeneratorPluginCreator::getPluginNamespace() const
TRT_NOEXCEPT {
return namespace_.c_str(); return namespace_.c_str();
} }
const char* AnchorGeneratorPluginCreator::getPluginName() const { const char* AnchorGeneratorPluginCreator::getPluginName() const TRT_NOEXCEPT {
return "anchor_generator_plugin"; return "anchor_generator_plugin";
} }
const char* AnchorGeneratorPluginCreator::getPluginVersion() const { const char* AnchorGeneratorPluginCreator::getPluginVersion() const
TRT_NOEXCEPT {
return "1"; return "1";
} }
const nvinfer1::PluginFieldCollection* const nvinfer1::PluginFieldCollection*
AnchorGeneratorPluginCreator::getFieldNames() { AnchorGeneratorPluginCreator::getFieldNames() TRT_NOEXCEPT {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2Ext* AnchorGeneratorPluginCreator::createPlugin( nvinfer1::IPluginV2Ext* AnchorGeneratorPluginCreator::createPlugin(
const char* name, const nvinfer1::PluginFieldCollection* fc) { const char* name, const nvinfer1::PluginFieldCollection* fc) TRT_NOEXCEPT {
const nvinfer1::PluginField* fields = fc->fields; const nvinfer1::PluginField* fields = fc->fields;
int type_id = -1; int type_id = -1;
std::vector<float> anchor_sizes, aspect_ratios, stride, variances; std::vector<float> anchor_sizes, aspect_ratios, stride, variances;
...@@ -315,7 +323,8 @@ nvinfer1::IPluginV2Ext* AnchorGeneratorPluginCreator::createPlugin( ...@@ -315,7 +323,8 @@ nvinfer1::IPluginV2Ext* AnchorGeneratorPluginCreator::createPlugin(
} }
nvinfer1::IPluginV2Ext* AnchorGeneratorPluginCreator::deserializePlugin( nvinfer1::IPluginV2Ext* AnchorGeneratorPluginCreator::deserializePlugin(
const char* name, const void* serial_data, size_t serial_length) { const char* name, const void* serial_data,
size_t serial_length) TRT_NOEXCEPT {
auto plugin = new AnchorGeneratorPlugin(serial_data, serial_length); auto plugin = new AnchorGeneratorPlugin(serial_data, serial_length);
plugin->setPluginNamespace(namespace_.c_str()); plugin->setPluginNamespace(namespace_.c_str());
return plugin; return plugin;
...@@ -374,7 +383,8 @@ AnchorGeneratorPluginDynamic::AnchorGeneratorPluginDynamic(void const* data, ...@@ -374,7 +383,8 @@ AnchorGeneratorPluginDynamic::AnchorGeneratorPluginDynamic(void const* data,
PrepareParamsOnDevice(); PrepareParamsOnDevice();
} }
nvinfer1::IPluginV2DynamicExt* AnchorGeneratorPluginDynamic::clone() const { nvinfer1::IPluginV2DynamicExt* AnchorGeneratorPluginDynamic::clone() const
TRT_NOEXCEPT {
auto plugin = new AnchorGeneratorPluginDynamic( auto plugin = new AnchorGeneratorPluginDynamic(
data_type_, anchor_sizes_, aspect_ratios_, stride_, variances_, offset_, data_type_, anchor_sizes_, aspect_ratios_, stride_, variances_, offset_,
num_anchors_); num_anchors_);
...@@ -384,7 +394,7 @@ nvinfer1::IPluginV2DynamicExt* AnchorGeneratorPluginDynamic::clone() const { ...@@ -384,7 +394,7 @@ nvinfer1::IPluginV2DynamicExt* AnchorGeneratorPluginDynamic::clone() const {
nvinfer1::DimsExprs AnchorGeneratorPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs AnchorGeneratorPluginDynamic::getOutputDimensions(
int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs, int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs,
nvinfer1::IExprBuilder& exprBuilder) { nvinfer1::IExprBuilder& exprBuilder) TRT_NOEXCEPT {
nvinfer1::DimsExprs ret{}; nvinfer1::DimsExprs ret{};
ret.nbDims = 4; ret.nbDims = 4;
ret.d[0] = inputs[0].d[2]; // feature height ret.d[0] = inputs[0].d[2]; // feature height
...@@ -396,7 +406,7 @@ nvinfer1::DimsExprs AnchorGeneratorPluginDynamic::getOutputDimensions( ...@@ -396,7 +406,7 @@ nvinfer1::DimsExprs AnchorGeneratorPluginDynamic::getOutputDimensions(
bool AnchorGeneratorPluginDynamic::supportsFormatCombination( bool AnchorGeneratorPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc* inOut, int nbInputs, int pos, const nvinfer1::PluginTensorDesc* inOut, int nbInputs,
int nbOutputs) { int nbOutputs) TRT_NOEXCEPT {
// input can be any, doesn't matter // input can be any, doesn't matter
// anchor generator doesn't read input raw data, only need the shape info // anchor generator doesn't read input raw data, only need the shape info
auto type = inOut[pos].type; auto type = inOut[pos].type;
...@@ -412,11 +422,12 @@ bool AnchorGeneratorPluginDynamic::supportsFormatCombination( ...@@ -412,11 +422,12 @@ bool AnchorGeneratorPluginDynamic::supportsFormatCombination(
void AnchorGeneratorPluginDynamic::configurePlugin( void AnchorGeneratorPluginDynamic::configurePlugin(
const nvinfer1::DynamicPluginTensorDesc* in, int nbInputs, const nvinfer1::DynamicPluginTensorDesc* in, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, int nbOutputs) {} const nvinfer1::DynamicPluginTensorDesc* out, int nbOutputs) TRT_NOEXCEPT {}
size_t AnchorGeneratorPluginDynamic::getWorkspaceSize( size_t AnchorGeneratorPluginDynamic::getWorkspaceSize(
const nvinfer1::PluginTensorDesc* inputs, int nbInputs, const nvinfer1::PluginTensorDesc* inputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, int nbOutputs) const { const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const TRT_NOEXCEPT {
return 0; return 0;
} }
...@@ -449,7 +460,7 @@ int AnchorGeneratorPluginDynamic::enqueue_impl( ...@@ -449,7 +460,7 @@ int AnchorGeneratorPluginDynamic::enqueue_impl(
int AnchorGeneratorPluginDynamic::enqueue( int AnchorGeneratorPluginDynamic::enqueue(
const nvinfer1::PluginTensorDesc* inputDesc, const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const void* const* inputs, const nvinfer1::PluginTensorDesc* outputDesc, const void* const* inputs,
void* const* outputs, void* workspace, cudaStream_t stream) { void* const* outputs, void* workspace, cudaStream_t stream) TRT_NOEXCEPT {
assert(outputDesc[0].type == nvinfer1::DataType::kFLOAT); assert(outputDesc[0].type == nvinfer1::DataType::kFLOAT);
assert(outputDesc[1].type == nvinfer1::DataType::kFLOAT); assert(outputDesc[1].type == nvinfer1::DataType::kFLOAT);
return enqueue_impl<float>(inputDesc, outputDesc, inputs, outputs, workspace, return enqueue_impl<float>(inputDesc, outputDesc, inputs, outputs, workspace,
...@@ -457,21 +468,24 @@ int AnchorGeneratorPluginDynamic::enqueue( ...@@ -457,21 +468,24 @@ int AnchorGeneratorPluginDynamic::enqueue(
} }
nvinfer1::DataType AnchorGeneratorPluginDynamic::getOutputDataType( nvinfer1::DataType AnchorGeneratorPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType* inputTypes, int nbInputs) const { int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const TRT_NOEXCEPT {
return inputTypes[0]; return inputTypes[0];
} }
const char* AnchorGeneratorPluginDynamic::getPluginType() const { const char* AnchorGeneratorPluginDynamic::getPluginType() const TRT_NOEXCEPT {
return "anchor_generator_plugin_dynamic"; return "anchor_generator_plugin_dynamic";
} }
int AnchorGeneratorPluginDynamic::getNbOutputs() const { return 2; } int AnchorGeneratorPluginDynamic::getNbOutputs() const TRT_NOEXCEPT {
return 2;
}
int AnchorGeneratorPluginDynamic::initialize() { return 0; } int AnchorGeneratorPluginDynamic::initialize() TRT_NOEXCEPT { return 0; }
void AnchorGeneratorPluginDynamic::terminate() {} void AnchorGeneratorPluginDynamic::terminate() TRT_NOEXCEPT {}
size_t AnchorGeneratorPluginDynamic::getSerializationSize() const { size_t AnchorGeneratorPluginDynamic::getSerializationSize() const TRT_NOEXCEPT {
size_t serialize_size = 0; size_t serialize_size = 0;
serialize_size += SerializedSize(data_type_); serialize_size += SerializedSize(data_type_);
serialize_size += SerializedSize(anchor_sizes_); serialize_size += SerializedSize(anchor_sizes_);
...@@ -483,7 +497,7 @@ size_t AnchorGeneratorPluginDynamic::getSerializationSize() const { ...@@ -483,7 +497,7 @@ size_t AnchorGeneratorPluginDynamic::getSerializationSize() const {
return serialize_size; return serialize_size;
} }
void AnchorGeneratorPluginDynamic::serialize(void* buffer) const { void AnchorGeneratorPluginDynamic::serialize(void* buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, data_type_); SerializeValue(&buffer, data_type_);
SerializeValue(&buffer, anchor_sizes_); SerializeValue(&buffer, anchor_sizes_);
SerializeValue(&buffer, aspect_ratios_); SerializeValue(&buffer, aspect_ratios_);
...@@ -493,32 +507,35 @@ void AnchorGeneratorPluginDynamic::serialize(void* buffer) const { ...@@ -493,32 +507,35 @@ void AnchorGeneratorPluginDynamic::serialize(void* buffer) const {
SerializeValue(&buffer, num_anchors_); SerializeValue(&buffer, num_anchors_);
} }
void AnchorGeneratorPluginDynamic::destroy() {} void AnchorGeneratorPluginDynamic::destroy() TRT_NOEXCEPT {}
void AnchorGeneratorPluginDynamicCreator::setPluginNamespace( void AnchorGeneratorPluginDynamicCreator::setPluginNamespace(
const char* lib_namespace) { const char* lib_namespace) TRT_NOEXCEPT {
namespace_ = std::string(lib_namespace); namespace_ = std::string(lib_namespace);
} }
const char* AnchorGeneratorPluginDynamicCreator::getPluginNamespace() const { const char* AnchorGeneratorPluginDynamicCreator::getPluginNamespace() const
TRT_NOEXCEPT {
return namespace_.c_str(); return namespace_.c_str();
} }
const char* AnchorGeneratorPluginDynamicCreator::getPluginName() const { const char* AnchorGeneratorPluginDynamicCreator::getPluginName() const
TRT_NOEXCEPT {
return "anchor_generator_plugin_dynamic"; return "anchor_generator_plugin_dynamic";
} }
const char* AnchorGeneratorPluginDynamicCreator::getPluginVersion() const { const char* AnchorGeneratorPluginDynamicCreator::getPluginVersion() const
TRT_NOEXCEPT {
return "1"; return "1";
} }
const nvinfer1::PluginFieldCollection* const nvinfer1::PluginFieldCollection*
AnchorGeneratorPluginDynamicCreator::getFieldNames() { AnchorGeneratorPluginDynamicCreator::getFieldNames() TRT_NOEXCEPT {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2Ext* AnchorGeneratorPluginDynamicCreator::createPlugin( nvinfer1::IPluginV2Ext* AnchorGeneratorPluginDynamicCreator::createPlugin(
const char* name, const nvinfer1::PluginFieldCollection* fc) { const char* name, const nvinfer1::PluginFieldCollection* fc) TRT_NOEXCEPT {
const nvinfer1::PluginField* fields = fc->fields; const nvinfer1::PluginField* fields = fc->fields;
int type_id = -1; int type_id = -1;
std::vector<float> anchor_sizes, aspect_ratios, stride, variances; std::vector<float> anchor_sizes, aspect_ratios, stride, variances;
...@@ -555,7 +572,8 @@ nvinfer1::IPluginV2Ext* AnchorGeneratorPluginDynamicCreator::createPlugin( ...@@ -555,7 +572,8 @@ nvinfer1::IPluginV2Ext* AnchorGeneratorPluginDynamicCreator::createPlugin(
} }
nvinfer1::IPluginV2Ext* AnchorGeneratorPluginDynamicCreator::deserializePlugin( nvinfer1::IPluginV2Ext* AnchorGeneratorPluginDynamicCreator::deserializePlugin(
const char* name, const void* serial_data, size_t serial_length) { const char* name, const void* serial_data,
size_t serial_length) TRT_NOEXCEPT {
auto plugin = new AnchorGeneratorPluginDynamic(serial_data, serial_length); auto plugin = new AnchorGeneratorPluginDynamic(serial_data, serial_length);
plugin->setPluginNamespace(namespace_.c_str()); plugin->setPluginNamespace(namespace_.c_str());
return plugin; return plugin;
......
paddle/fluid/inference/tensorrt/plugin/anchor_generator_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -34,34 +34,35 @@ class AnchorGeneratorPlugin : public nvinfer1::IPluginV2Ext { ...@@ -34,34 +34,35 @@ class AnchorGeneratorPlugin : public nvinfer1::IPluginV2Ext {
const int width, const int num_anchors, const int box_num); const int width, const int num_anchors, const int box_num);
AnchorGeneratorPlugin(const void* data, size_t length); AnchorGeneratorPlugin(const void* data, size_t length);
~AnchorGeneratorPlugin() override; ~AnchorGeneratorPlugin() override;
const char* getPluginType() const override; const char* getPluginType() const TRT_NOEXCEPT override;
const char* getPluginVersion() const override; const char* getPluginVersion() const TRT_NOEXCEPT override;
int getNbOutputs() const override; int getNbOutputs() const TRT_NOEXCEPT override;
nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs,
int nb_input_dims) override; int nb_input_dims) TRT_NOEXCEPT override;
bool supportsFormat(nvinfer1::DataType type, bool supportsFormat(nvinfer1::DataType type, nvinfer1::TensorFormat format)
nvinfer1::TensorFormat format) const override; const TRT_NOEXCEPT override;
size_t getWorkspaceSize(int max_batch_size) const override; size_t getWorkspaceSize(int max_batch_size) const TRT_NOEXCEPT override;
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
int enqueue(int batch_size, const void* const* inputs, void** outputs, int enqueue(int batch_size, const void* const* inputs, void** outputs,
#else #else
int enqueue(int batch_size, const void* const* inputs, void* const* outputs, int enqueue(int batch_size, const void* const* inputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream) override; void* workspace, cudaStream_t stream) TRT_NOEXCEPT override;
int initialize() override; int initialize() TRT_NOEXCEPT override;
void terminate() override; void terminate() TRT_NOEXCEPT override;
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
void destroy() override; void destroy() TRT_NOEXCEPT override;
void setPluginNamespace(const char* lib_namespace) override; void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override;
const char* getPluginNamespace() const override; const char* getPluginNamespace() const TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* input_type, int index, const nvinfer1::DataType* input_type,
int nb_inputs) const override; int nb_inputs) const TRT_NOEXCEPT override;
bool isOutputBroadcastAcrossBatch(int output_index, bool isOutputBroadcastAcrossBatch(int output_index,
const bool* input_is_broadcast, const bool* input_is_broadcast,
int nb_inputs) const override; int nb_inputs) const TRT_NOEXCEPT override;
bool canBroadcastInputAcrossBatch(int input_index) const override; bool canBroadcastInputAcrossBatch(int input_index) const
TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::Dims* input_dims, int nb_inputs, void configurePlugin(const nvinfer1::Dims* input_dims, int nb_inputs,
const nvinfer1::Dims* output_dims, int nb_outputs, const nvinfer1::Dims* output_dims, int nb_outputs,
const nvinfer1::DataType* input_types, const nvinfer1::DataType* input_types,
...@@ -69,13 +70,13 @@ class AnchorGeneratorPlugin : public nvinfer1::IPluginV2Ext { ...@@ -69,13 +70,13 @@ class AnchorGeneratorPlugin : public nvinfer1::IPluginV2Ext {
const bool* input_is_broadcast, const bool* input_is_broadcast,
const bool* output_is_broadcast, const bool* output_is_broadcast,
nvinfer1::PluginFormat float_format, nvinfer1::PluginFormat float_format,
int max_batct_size) override; int max_batct_size) TRT_NOEXCEPT override;
nvinfer1::IPluginV2Ext* clone() const override; nvinfer1::IPluginV2Ext* clone() const TRT_NOEXCEPT override;
private: private:
template <typename T> template <typename T>
int enqueue_impl(int batch_size, const void* const* inputs, void** outputs, int enqueue_impl(int batch_size, const void* const* inputs,
void* workspace, cudaStream_t stream); void* const* outputs, void* workspace, cudaStream_t stream);
nvinfer1::DataType data_type_; nvinfer1::DataType data_type_;
std::vector<float> anchor_sizes_; std::vector<float> anchor_sizes_;
std::vector<float> aspect_ratios_; std::vector<float> aspect_ratios_;
...@@ -97,16 +98,17 @@ class AnchorGeneratorPluginCreator : public nvinfer1::IPluginCreator { ...@@ -97,16 +98,17 @@ class AnchorGeneratorPluginCreator : public nvinfer1::IPluginCreator {
public: public:
AnchorGeneratorPluginCreator() = default; AnchorGeneratorPluginCreator() = default;
~AnchorGeneratorPluginCreator() override = default; ~AnchorGeneratorPluginCreator() override = default;
void setPluginNamespace(const char* lib_namespace) override; void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override;
const char* getPluginNamespace() const override; const char* getPluginNamespace() const TRT_NOEXCEPT override;
const char* getPluginName() const override; const char* getPluginName() const TRT_NOEXCEPT override;
const char* getPluginVersion() const override; const char* getPluginVersion() const TRT_NOEXCEPT override;
const nvinfer1::PluginFieldCollection* getFieldNames() override; const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override;
nvinfer1::IPluginV2Ext* createPlugin( nvinfer1::IPluginV2Ext* createPlugin(
const char* name, const nvinfer1::PluginFieldCollection* fc) override; const char* name,
nvinfer1::IPluginV2Ext* deserializePlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) TRT_NOEXCEPT override;
const void* serial_data, nvinfer1::IPluginV2Ext* deserializePlugin(
size_t serial_length) override; const char* name, const void* serial_data,
size_t serial_length) TRT_NOEXCEPT override;
private: private:
std::string namespace_; std::string namespace_;
...@@ -127,35 +129,36 @@ class AnchorGeneratorPluginDynamic : public DynamicPluginTensorRT { ...@@ -127,35 +129,36 @@ class AnchorGeneratorPluginDynamic : public DynamicPluginTensorRT {
const int num_anchors); const int num_anchors);
AnchorGeneratorPluginDynamic(void const* data, size_t length); AnchorGeneratorPluginDynamic(void const* data, size_t length);
~AnchorGeneratorPluginDynamic(); ~AnchorGeneratorPluginDynamic();
nvinfer1::IPluginV2DynamicExt* clone() const override; nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override;
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs, int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs,
nvinfer1::IExprBuilder& exprBuilder) override; nvinfer1::IExprBuilder& exprBuilder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* inOut, const nvinfer1::PluginTensorDesc* inOut,
int nbInputs, int nbOutputs) override; int nbInputs,
int nbOutputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nbInputs, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nbOutputs) override; int nbOutputs) TRT_NOEXCEPT override;
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nbInputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const override; int nbOutputs) const TRT_NOEXCEPT override;
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* inputTypes, int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const override; int nbInputs) const TRT_NOEXCEPT override;
const char* getPluginType() const override; const char* getPluginType() const TRT_NOEXCEPT override;
int getNbOutputs() const override; int getNbOutputs() const TRT_NOEXCEPT override;
int initialize() override; int initialize() TRT_NOEXCEPT override;
void terminate() override; void terminate() TRT_NOEXCEPT override;
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
void destroy() override; void destroy() TRT_NOEXCEPT override;
private: private:
template <typename T> template <typename T>
...@@ -181,16 +184,17 @@ class AnchorGeneratorPluginDynamicCreator : public nvinfer1::IPluginCreator { ...@@ -181,16 +184,17 @@ class AnchorGeneratorPluginDynamicCreator : public nvinfer1::IPluginCreator {
public: public:
AnchorGeneratorPluginDynamicCreator() = default; AnchorGeneratorPluginDynamicCreator() = default;
~AnchorGeneratorPluginDynamicCreator() override = default; ~AnchorGeneratorPluginDynamicCreator() override = default;
void setPluginNamespace(const char* lib_namespace) override; void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override;
const char* getPluginNamespace() const override; const char* getPluginNamespace() const TRT_NOEXCEPT override;
const char* getPluginName() const override; const char* getPluginName() const TRT_NOEXCEPT override;
const char* getPluginVersion() const override; const char* getPluginVersion() const TRT_NOEXCEPT override;
const nvinfer1::PluginFieldCollection* getFieldNames() override; const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override;
nvinfer1::IPluginV2Ext* createPlugin( nvinfer1::IPluginV2Ext* createPlugin(
const char* name, const nvinfer1::PluginFieldCollection* fc) override; const char* name,
nvinfer1::IPluginV2Ext* deserializePlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) TRT_NOEXCEPT override;
const void* serial_data, nvinfer1::IPluginV2Ext* deserializePlugin(
size_t serial_length) override; const char* name, const void* serial_data,
size_t serial_length) TRT_NOEXCEPT override;
private: private:
std::string namespace_; std::string namespace_;
......
paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -48,7 +48,7 @@ __global__ void elementwise_kernel(const size_t total, const T *x_data, ...@@ -48,7 +48,7 @@ __global__ void elementwise_kernel(const size_t total, const T *x_data,
} }
nvinfer1::Dims ElementWisePlugin::getOutputDimensions( nvinfer1::Dims ElementWisePlugin::getOutputDimensions(
int index, const nvinfer1::Dims *input_dims, int num_inputs) { int index, const nvinfer1::Dims *input_dims, int num_inputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument( PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument(
"There is only one output in TRT elementwise " "There is only one output in TRT elementwise "
"op plugin, but got output index: %d.", "op plugin, but got output index: %d.",
...@@ -64,7 +64,7 @@ nvinfer1::Dims ElementWisePlugin::getOutputDimensions( ...@@ -64,7 +64,7 @@ nvinfer1::Dims ElementWisePlugin::getOutputDimensions(
return input_dims[0]; return input_dims[0];
} }
int ElementWisePlugin::initialize() { int ElementWisePlugin::initialize() TRT_NOEXCEPT {
PADDLE_ENFORCE_GT(dims_y_.nbDims, 0, PADDLE_ENFORCE_GT(dims_y_.nbDims, 0,
platform::errors::InvalidArgument( platform::errors::InvalidArgument(
"The dimension of input Y of TRT elementwise op plugin " "The dimension of input Y of TRT elementwise op plugin "
...@@ -120,7 +120,7 @@ int ElementWisePlugin::enqueue(int batch_size, const void *const *inputs, ...@@ -120,7 +120,7 @@ int ElementWisePlugin::enqueue(int batch_size, const void *const *inputs,
#else #else
void *const *outputs, void *workspace, void *const *outputs, void *workspace,
#endif #endif
cudaStream_t stream) { cudaStream_t stream) TRT_NOEXCEPT {
const float *x = reinterpret_cast<const float *>(inputs[0]); const float *x = reinterpret_cast<const float *>(inputs[0]);
const float *y = reinterpret_cast<const float *>(inputs[1]); const float *y = reinterpret_cast<const float *>(inputs[1]);
float *out = reinterpret_cast<float *>(outputs[0]); float *out = reinterpret_cast<float *>(outputs[0]);
...@@ -147,26 +147,26 @@ int ElementWisePlugin::enqueue(int batch_size, const void *const *inputs, ...@@ -147,26 +147,26 @@ int ElementWisePlugin::enqueue(int batch_size, const void *const *inputs,
// Dynamic Plugin below. // Dynamic Plugin below.
#if IS_TRT_VERSION_GE(6000) #if IS_TRT_VERSION_GE(6000)
int ElementwisePluginDynamic::initialize() { return 0; } int ElementwisePluginDynamic::initialize() TRT_NOEXCEPT { return 0; }
size_t ElementwisePluginDynamic::getSerializationSize() const { size_t ElementwisePluginDynamic::getSerializationSize() const TRT_NOEXCEPT {
return SerializedSize(type_.c_str()) + SerializedSize(axis_); return SerializedSize(type_.c_str()) + SerializedSize(axis_);
} }
void ElementwisePluginDynamic::serialize(void *buffer) const { void ElementwisePluginDynamic::serialize(void *buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, type_.c_str()); SerializeValue(&buffer, type_.c_str());
SerializeValue(&buffer, axis_); SerializeValue(&buffer, axis_);
} }
nvinfer1::DimsExprs ElementwisePluginDynamic::getOutputDimensions( nvinfer1::DimsExprs ElementwisePluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs,
nvinfer1::IExprBuilder &expr_builder) { nvinfer1::IExprBuilder &expr_builder) TRT_NOEXCEPT {
return inputs[0]; return inputs[0];
} }
bool ElementwisePluginDynamic::supportsFormatCombination( bool ElementwisePluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of swish plugin shoule not be nullptr.")); "The input of swish plugin shoule not be nullptr."));
...@@ -189,7 +189,8 @@ bool ElementwisePluginDynamic::supportsFormatCombination( ...@@ -189,7 +189,8 @@ bool ElementwisePluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType ElementwisePluginDynamic::getOutputDataType( nvinfer1::DataType ElementwisePluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType *input_types, int nb_inputs) const { int index, const nvinfer1::DataType *input_types,
int nb_inputs) const TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(index, 0, PADDLE_ENFORCE_EQ(index, 0,
platform::errors::InvalidArgument( platform::errors::InvalidArgument(
"The Elementwise Plugin only has one input, so the " "The Elementwise Plugin only has one input, so the "
...@@ -201,7 +202,7 @@ nvinfer1::DataType ElementwisePluginDynamic::getOutputDataType( ...@@ -201,7 +202,7 @@ nvinfer1::DataType ElementwisePluginDynamic::getOutputDataType(
int ElementwisePluginDynamic::enqueue( int ElementwisePluginDynamic::enqueue(
const nvinfer1::PluginTensorDesc *input_desc, const nvinfer1::PluginTensorDesc *input_desc,
const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs, const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs,
void *const *outputs, void *workspace, cudaStream_t stream) { void *const *outputs, void *workspace, cudaStream_t stream) TRT_NOEXCEPT {
auto x_dims = input_desc[0].dims; auto x_dims = input_desc[0].dims;
auto y_dims = input_desc[1].dims; auto y_dims = input_desc[1].dims;
int axis = (axis_ == -1) ? x_dims.nbDims - y_dims.nbDims : axis_; int axis = (axis_ == -1) ? x_dims.nbDims - y_dims.nbDims : axis_;
......
paddle/fluid/inference/tensorrt/plugin/elementwise_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -48,33 +48,35 @@ class ElementWisePlugin : public PluginTensorRT { ...@@ -48,33 +48,35 @@ class ElementWisePlugin : public PluginTensorRT {
DeserializeValue(&serial_data, &serial_length, &post_size_); DeserializeValue(&serial_data, &serial_length, &post_size_);
} }
ElementWisePlugin* clone() const override { ElementWisePlugin* clone() const TRT_NOEXCEPT override {
return new ElementWisePlugin(type_, dims_x_, dims_y_, axis_); return new ElementWisePlugin(type_, dims_x_, dims_y_, axis_);
} }
const char* getPluginType() const override { return "elementwise_plugin"; } const char* getPluginType() const TRT_NOEXCEPT override {
return "elementwise_plugin";
}
nvinfer1::Dims getOutputDimensions(int index, nvinfer1::Dims getOutputDimensions(int index,
const nvinfer1::Dims* input_dims, const nvinfer1::Dims* input_dims,
int num_inputs) override; int num_inputs) TRT_NOEXCEPT override;
int initialize() override; int initialize() TRT_NOEXCEPT override;
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
int enqueue(int batch_size, const void* const* inputs, void** outputs, int enqueue(int batch_size, const void* const* inputs, void** outputs,
#else #else
int enqueue(int batch_size, const void* const* inputs, void* const* outputs, int enqueue(int batch_size, const void* const* inputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream); void* workspace, cudaStream_t stream) TRT_NOEXCEPT;
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
return getBaseSerializationSize() + SerializedSize(type_.c_str()) + return getBaseSerializationSize() + SerializedSize(type_.c_str()) +
SerializedSize(dims_x_) + SerializedSize(dims_y_) + SerializedSize(dims_x_) + SerializedSize(dims_y_) +
SerializedSize(axis_) + SerializedSize(prev_size_) + SerializedSize(axis_) + SerializedSize(prev_size_) +
SerializedSize(midd_size_) + SerializedSize(post_size_); SerializedSize(midd_size_) + SerializedSize(post_size_);
} }
void serialize(void* buffer) const override { void serialize(void* buffer) const TRT_NOEXCEPT override {
serializeBase(buffer); serializeBase(buffer);
SerializeValue(&buffer, type_.c_str()); SerializeValue(&buffer, type_.c_str());
SerializeValue(&buffer, dims_x_); SerializeValue(&buffer, dims_x_);
...@@ -97,13 +99,15 @@ class ElementWisePlugin : public PluginTensorRT { ...@@ -97,13 +99,15 @@ class ElementWisePlugin : public PluginTensorRT {
class ElementWisePluginCreator : public TensorRTPluginCreator { class ElementWisePluginCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "elementwise_plugin"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "elementwise_plugin";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new ElementWisePlugin(serial_data, serial_length); return new ElementWisePlugin(serial_data, serial_length);
} }
}; };
...@@ -120,48 +124,49 @@ class ElementwisePluginDynamic : public DynamicPluginTensorRT { ...@@ -120,48 +124,49 @@ class ElementwisePluginDynamic : public DynamicPluginTensorRT {
type_ = std::string(elementwise_type); type_ = std::string(elementwise_type);
DeserializeValue(&serialData, &serialLength, &axis_); DeserializeValue(&serialData, &serialLength, &axis_);
} }
nvinfer1::IPluginV2DynamicExt* clone() const override { nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override {
return new ElementwisePluginDynamic(type_, axis_); return new ElementwisePluginDynamic(type_, axis_);
} }
const char* getPluginType() const override { const char* getPluginType() const TRT_NOEXCEPT override {
return "elementwise_plugin_dynamic"; return "elementwise_plugin_dynamic";
} }
int getNbOutputs() const override { return 1; } int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
int initialize() override; int initialize() TRT_NOEXCEPT override;
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) override; nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* inOut, const nvinfer1::PluginTensorDesc* inOut,
int nbInputs, int nbOutputs) override; int nbInputs,
int nbOutputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nbInputs, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nbOutputs) override {} int nbOutputs) TRT_NOEXCEPT override {}
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nbInputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const override { int nbOutputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* inputTypes, int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const override; int nbInputs) const TRT_NOEXCEPT override;
void destroy() override { delete this; } void destroy() TRT_NOEXCEPT override { delete this; }
private: private:
std::string type_; std::string type_;
...@@ -171,33 +176,34 @@ class ElementwisePluginDynamic : public DynamicPluginTensorRT { ...@@ -171,33 +176,34 @@ class ElementwisePluginDynamic : public DynamicPluginTensorRT {
class ElementwisePluginDynamicCreator : public nvinfer1::IPluginCreator { class ElementwisePluginDynamicCreator : public nvinfer1::IPluginCreator {
public: public:
ElementwisePluginDynamicCreator() {} ElementwisePluginDynamicCreator() {}
const char* getPluginName() const override { const char* getPluginName() const TRT_NOEXCEPT override {
return "elementwise_plugin_dynamic"; return "elementwise_plugin_dynamic";
} }
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
const nvinfer1::PluginFieldCollection* getFieldNames() override { const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2* createPlugin( nvinfer1::IPluginV2* createPlugin(const char* name,
const char* name, const nvinfer1::PluginFieldCollection* fc) override { const nvinfer1::PluginFieldCollection* fc)
TRT_NOEXCEPT override {
return nullptr; return nullptr;
} }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
auto plugin = new ElementwisePluginDynamic(serial_data, serial_length); auto plugin = new ElementwisePluginDynamic(serial_data, serial_length);
return plugin; return plugin;
} }
void setPluginNamespace(const char* lib_namespace) override { void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override {
plugin_namespace_ = lib_namespace; plugin_namespace_ = lib_namespace;
} }
const char* getPluginNamespace() const override { const char* getPluginNamespace() const TRT_NOEXCEPT override {
return plugin_namespace_.c_str(); return plugin_namespace_.c_str();
} }
......
paddle/fluid/inference/tensorrt/plugin/emb_eltwise_layernorm_plugin.cu
浏览文件 @ b5aab4f0
...@@ -128,7 +128,7 @@ template <typename T> ...@@ -128,7 +128,7 @@ template <typename T>
int EmbEltwiseLayernormPluginDynamicImpl<T>::enqueue( int EmbEltwiseLayernormPluginDynamicImpl<T>::enqueue(
const nvinfer1::PluginTensorDesc *input_desc, const nvinfer1::PluginTensorDesc *input_desc,
const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs, const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs,
void *const *outputs, void *workspace, cudaStream_t stream) { void *const *outputs, void *workspace, cudaStream_t stream) TRT_NOEXCEPT {
auto id_dims = input_desc[0].dims; auto id_dims = input_desc[0].dims;
int batch = id_dims.d[0]; int batch = id_dims.d[0];
int seq_len = id_dims.d[1]; int seq_len = id_dims.d[1];
...@@ -181,17 +181,19 @@ template class EmbEltwiseLayernormPluginDynamicImpl<float>; ...@@ -181,17 +181,19 @@ template class EmbEltwiseLayernormPluginDynamicImpl<float>;
template class EmbEltwiseLayernormPluginDynamicImpl<half>; template class EmbEltwiseLayernormPluginDynamicImpl<half>;
#endif #endif
int EmbEltwiseLayernormPluginDynamic::initialize() { int EmbEltwiseLayernormPluginDynamic::initialize() TRT_NOEXCEPT {
impl_->initialize(); impl_->initialize();
return 0; return 0;
} }
void EmbEltwiseLayernormPluginDynamic::terminate() { impl_->terminate(); } void EmbEltwiseLayernormPluginDynamic::terminate() TRT_NOEXCEPT {
impl_->terminate();
}
nvinfer1::DimsExprs EmbEltwiseLayernormPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs EmbEltwiseLayernormPluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs,
nvinfer1::IExprBuilder &expr_builder) { // NOLINT nvinfer1::IExprBuilder &expr_builder) TRT_NOEXCEPT { // NOLINT
PADDLE_ENFORCE_EQ(output_index, 0, PADDLE_ENFORCE_EQ(output_index, 0,
platform::errors::InvalidArgument( platform::errors::InvalidArgument(
"There is only one output of the EmbEltwiseLayernorm, " "There is only one output of the EmbEltwiseLayernorm, "
...@@ -208,7 +210,7 @@ nvinfer1::DimsExprs EmbEltwiseLayernormPluginDynamic::getOutputDimensions( ...@@ -208,7 +210,7 @@ nvinfer1::DimsExprs EmbEltwiseLayernormPluginDynamic::getOutputDimensions(
bool EmbEltwiseLayernormPluginDynamic::supportsFormatCombination( bool EmbEltwiseLayernormPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of swish plugin shoule not be nullptr.")); "The input of swish plugin shoule not be nullptr."));
...@@ -256,7 +258,8 @@ bool EmbEltwiseLayernormPluginDynamic::supportsFormatCombination( ...@@ -256,7 +258,8 @@ bool EmbEltwiseLayernormPluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType EmbEltwiseLayernormPluginDynamic::getOutputDataType( nvinfer1::DataType EmbEltwiseLayernormPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType *input_types, int nb_inputs) const { int index, const nvinfer1::DataType *input_types,
int nb_inputs) const TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ( PADDLE_ENFORCE_EQ(
index, 0, platform::errors::InvalidArgument( index, 0, platform::errors::InvalidArgument(
"The EmbEltwiseLayernorm Plugin only has one input, so the " "The EmbEltwiseLayernorm Plugin only has one input, so the "
...@@ -271,7 +274,7 @@ nvinfer1::DataType EmbEltwiseLayernormPluginDynamic::getOutputDataType( ...@@ -271,7 +274,7 @@ nvinfer1::DataType EmbEltwiseLayernormPluginDynamic::getOutputDataType(
int EmbEltwiseLayernormPluginDynamic::enqueue( int EmbEltwiseLayernormPluginDynamic::enqueue(
const nvinfer1::PluginTensorDesc *input_desc, const nvinfer1::PluginTensorDesc *input_desc,
const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs, const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs,
void *const *outputs, void *workspace, cudaStream_t stream) { void *const *outputs, void *workspace, cudaStream_t stream) TRT_NOEXCEPT {
impl_->enqueue(input_desc, output_desc, inputs, outputs, workspace, stream); impl_->enqueue(input_desc, output_desc, inputs, outputs, workspace, stream);
return cudaGetLastError() != cudaSuccess; return cudaGetLastError() != cudaSuccess;
} }
......
paddle/fluid/inference/tensorrt/plugin/emb_eltwise_layernorm_plugin.h
浏览文件 @ b5aab4f0
...@@ -68,7 +68,7 @@ class EmbEltwiseLayernormPluginDynamicImpl ...@@ -68,7 +68,7 @@ class EmbEltwiseLayernormPluginDynamicImpl
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream); cudaStream_t stream) TRT_NOEXCEPT;
void shareGPUData(const EmbEltwiseLayernormPluginDynamicImplBase* anthor); void shareGPUData(const EmbEltwiseLayernormPluginDynamicImplBase* anthor);
private: private:
...@@ -189,7 +189,7 @@ class EmbEltwiseLayernormPluginDynamic : public DynamicPluginTensorRT { ...@@ -189,7 +189,7 @@ class EmbEltwiseLayernormPluginDynamic : public DynamicPluginTensorRT {
} }
} }
nvinfer1::IPluginV2DynamicExt* clone() const override { nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override {
auto ptr = new EmbEltwiseLayernormPluginDynamic( auto ptr = new EmbEltwiseLayernormPluginDynamic(
embs_, bias_, scale_, emb_sizes_, bias_size_, scale_size_, hidden_size_, embs_, bias_, scale_, emb_sizes_, bias_size_, scale_size_, hidden_size_,
eps_, with_fp16_); eps_, with_fp16_);
...@@ -197,14 +197,14 @@ class EmbEltwiseLayernormPluginDynamic : public DynamicPluginTensorRT { ...@@ -197,14 +197,14 @@ class EmbEltwiseLayernormPluginDynamic : public DynamicPluginTensorRT {
return ptr; return ptr;
} }
const char* getPluginType() const override { const char* getPluginType() const TRT_NOEXCEPT override {
return "fused_embedding_eltwise_layernorm_plugin"; return "fused_embedding_eltwise_layernorm_plugin";
} }
int getNbOutputs() const override { return 1; } int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
int initialize() override; int initialize() TRT_NOEXCEPT override;
void terminate() override; void terminate() TRT_NOEXCEPT override;
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
int sum_num = 0; int sum_num = 0;
sum_num += SerializedSize(emb_sizes_); sum_num += SerializedSize(emb_sizes_);
...@@ -223,7 +223,7 @@ class EmbEltwiseLayernormPluginDynamic : public DynamicPluginTensorRT { ...@@ -223,7 +223,7 @@ class EmbEltwiseLayernormPluginDynamic : public DynamicPluginTensorRT {
return sum_num; return sum_num;
} }
void serialize(void* buffer) const override { void serialize(void* buffer) const TRT_NOEXCEPT override {
SerializeValue(&buffer, emb_sizes_); SerializeValue(&buffer, emb_sizes_);
for (size_t i = 0; i < emb_sizes_.size(); i++) { for (size_t i = 0; i < emb_sizes_.size(); i++) {
auto size = emb_sizes_[i]; auto size = emb_sizes_[i];
...@@ -248,33 +248,34 @@ class EmbEltwiseLayernormPluginDynamic : public DynamicPluginTensorRT { ...@@ -248,33 +248,34 @@ class EmbEltwiseLayernormPluginDynamic : public DynamicPluginTensorRT {
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) override; nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* in_out, const nvinfer1::PluginTensorDesc* in_out,
int nb_inputs, int nb_outputs) override; int nb_inputs,
int nb_outputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nb_inputs, int nb_inputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nb_outputs) override {} int nb_outputs) TRT_NOEXCEPT override {}
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nb_inputs, int nb_inputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nb_outputs) const override { int nb_outputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
int enqueue(const nvinfer1::PluginTensorDesc* input_desc, int enqueue(const nvinfer1::PluginTensorDesc* input_desc,
const nvinfer1::PluginTensorDesc* output_desc, const nvinfer1::PluginTensorDesc* output_desc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* input_types, int index, const nvinfer1::DataType* input_types,
int nb_inputs) const override; int nb_inputs) const TRT_NOEXCEPT override;
void destroy() override { void destroy() TRT_NOEXCEPT override {
if (own_host_buff_) { if (own_host_buff_) {
for (auto ptr : embs_) { for (auto ptr : embs_) {
delete[] ptr; delete[] ptr;
...@@ -310,32 +311,33 @@ class EmbEltwiseLayernormPluginDynamicCreator ...@@ -310,32 +311,33 @@ class EmbEltwiseLayernormPluginDynamicCreator
: public nvinfer1::IPluginCreator { : public nvinfer1::IPluginCreator {
public: public:
EmbEltwiseLayernormPluginDynamicCreator() {} EmbEltwiseLayernormPluginDynamicCreator() {}
const char* getPluginName() const override { const char* getPluginName() const TRT_NOEXCEPT override {
return "fused_embedding_eltwise_layernorm_plugin"; return "fused_embedding_eltwise_layernorm_plugin";
} }
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
const nvinfer1::PluginFieldCollection* getFieldNames() override { const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2* createPlugin( nvinfer1::IPluginV2* createPlugin(const char* name,
const char* name, const nvinfer1::PluginFieldCollection* fc) override { const nvinfer1::PluginFieldCollection* fc)
TRT_NOEXCEPT override {
return nullptr; return nullptr;
} }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new EmbEltwiseLayernormPluginDynamic(serial_data, serial_length); return new EmbEltwiseLayernormPluginDynamic(serial_data, serial_length);
} }
void setPluginNamespace(const char* lib_namespace) override { void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override {
plugin_namespace_ = lib_namespace; plugin_namespace_ = lib_namespace;
} }
const char* getPluginNamespace() const override { const char* getPluginNamespace() const TRT_NOEXCEPT override {
return plugin_namespace_.c_str(); return plugin_namespace_.c_str();
} }
......
paddle/fluid/inference/tensorrt/plugin/gather_nd_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -58,19 +58,19 @@ __global__ void GatherNdCUDAKernel(const T* input, const int32_t* input_dims, ...@@ -58,19 +58,19 @@ __global__ void GatherNdCUDAKernel(const T* input, const int32_t* input_dims,
} }
} }
int GatherNdPluginDynamic::initialize() { return 0; } int GatherNdPluginDynamic::initialize() TRT_NOEXCEPT { return 0; }
size_t GatherNdPluginDynamic::getSerializationSize() const { size_t GatherNdPluginDynamic::getSerializationSize() const TRT_NOEXCEPT {
return SerializedSize(with_fp16_); return SerializedSize(with_fp16_);
} }
void GatherNdPluginDynamic::serialize(void* buffer) const { void GatherNdPluginDynamic::serialize(void* buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, with_fp16_); SerializeValue(&buffer, with_fp16_);
} }
nvinfer1::DimsExprs GatherNdPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs GatherNdPluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) { nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ( PADDLE_ENFORCE_EQ(
nb_inputs, 2, nb_inputs, 2,
platform::errors::InvalidArgument( platform::errors::InvalidArgument(
...@@ -100,7 +100,7 @@ nvinfer1::DimsExprs GatherNdPluginDynamic::getOutputDimensions( ...@@ -100,7 +100,7 @@ nvinfer1::DimsExprs GatherNdPluginDynamic::getOutputDimensions(
bool GatherNdPluginDynamic::supportsFormatCombination( bool GatherNdPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc* in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc* in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of gather_nd plugin should not be nullptr.")); "The input of gather_nd plugin should not be nullptr."));
...@@ -134,14 +134,15 @@ bool GatherNdPluginDynamic::supportsFormatCombination( ...@@ -134,14 +134,15 @@ bool GatherNdPluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType GatherNdPluginDynamic::getOutputDataType( nvinfer1::DataType GatherNdPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType* input_types, int nb_inputs) const { int index, const nvinfer1::DataType* input_types,
int nb_inputs) const TRT_NOEXCEPT {
return input_types[0]; return input_types[0];
} }
int GatherNdPluginDynamic::enqueue( int GatherNdPluginDynamic::enqueue(
const nvinfer1::PluginTensorDesc* input_desc, const nvinfer1::PluginTensorDesc* input_desc,
const nvinfer1::PluginTensorDesc* output_desc, const void* const* inputs, const nvinfer1::PluginTensorDesc* output_desc, const void* const* inputs,
void* const* outputs, void* workspace, cudaStream_t stream) { void* const* outputs, void* workspace, cudaStream_t stream) TRT_NOEXCEPT {
auto input_dims = input_desc[0].dims; auto input_dims = input_desc[0].dims;
auto index_dims = input_desc[1].dims; auto index_dims = input_desc[1].dims;
auto input_dims_size = input_dims.nbDims; auto input_dims_size = input_dims.nbDims;
......
paddle/fluid/inference/tensorrt/plugin/gather_nd_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -35,46 +35,49 @@ class GatherNdPluginDynamic : public DynamicPluginTensorRT { ...@@ -35,46 +35,49 @@ class GatherNdPluginDynamic : public DynamicPluginTensorRT {
DeserializeValue(&serial_data, &serial_length, &with_fp16_); DeserializeValue(&serial_data, &serial_length, &with_fp16_);
} }
nvinfer1::IPluginV2DynamicExt* clone() const override { nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override {
return new GatherNdPluginDynamic(with_fp16_); return new GatherNdPluginDynamic(with_fp16_);
} }
const char* getPluginType() const override { return "gather_nd_plugin"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "gather_nd_plugin";
int initialize() override; }
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
int initialize() TRT_NOEXCEPT override;
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs, int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs,
nvinfer1::IExprBuilder& exprBuilder) override; nvinfer1::IExprBuilder& exprBuilder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* inOut, const nvinfer1::PluginTensorDesc* inOut,
int nbInputs, int nbOutputs) override; int nbInputs,
int nbOutputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nbInputs, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nbOutputs) override {} int nbOutputs) TRT_NOEXCEPT override {}
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nbInputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const override { int nbOutputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* inputTypes, int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const override; int nbInputs) const TRT_NOEXCEPT override;
void destroy() override { void destroy() TRT_NOEXCEPT override {
if (input_dims_data_) { if (input_dims_data_) {
cudaFree(input_dims_data_); cudaFree(input_dims_data_);
} }
...@@ -88,31 +91,34 @@ class GatherNdPluginDynamic : public DynamicPluginTensorRT { ...@@ -88,31 +91,34 @@ class GatherNdPluginDynamic : public DynamicPluginTensorRT {
class GatherNdPluginDynamicCreator : public nvinfer1::IPluginCreator { class GatherNdPluginDynamicCreator : public nvinfer1::IPluginCreator {
public: public:
GatherNdPluginDynamicCreator() {} GatherNdPluginDynamicCreator() {}
const char* getPluginName() const override { return "gather_nd_plugin"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "gather_nd_plugin";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
const nvinfer1::PluginFieldCollection* getFieldNames() override { const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2* createPlugin( nvinfer1::IPluginV2* createPlugin(const char* name,
const char* name, const nvinfer1::PluginFieldCollection* fc) override { const nvinfer1::PluginFieldCollection* fc)
TRT_NOEXCEPT override {
return nullptr; return nullptr;
} }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
auto plugin = new GatherNdPluginDynamic(serial_data, serial_length); auto plugin = new GatherNdPluginDynamic(serial_data, serial_length);
return plugin; return plugin;
} }
void setPluginNamespace(const char* lib_namespace) override { void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override {
plugin_namespace_ = lib_namespace; plugin_namespace_ = lib_namespace;
} }
const char* getPluginNamespace() const override { const char* getPluginNamespace() const TRT_NOEXCEPT override {
return plugin_namespace_.c_str(); return plugin_namespace_.c_str();
} }
......
paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -30,8 +30,8 @@ static const float kAT = 0.5; ...@@ -30,8 +30,8 @@ static const float kAT = 0.5;
static const float kBT = 0.7978845608028654; // sqrt(2.0/M_PI) static const float kBT = 0.7978845608028654; // sqrt(2.0/M_PI)
static const float kCT = 0.035677408136300125; // 0.044715 * sqrt(2.0/M_PI) static const float kCT = 0.035677408136300125; // 0.044715 * sqrt(2.0/M_PI)
bool GeluPlugin::supportsFormat(nvinfer1::DataType type, bool GeluPlugin::supportsFormat(
nvinfer1::PluginFormat format) const { nvinfer1::DataType type, nvinfer1::PluginFormat format) const TRT_NOEXCEPT {
if (with_fp16_) { if (with_fp16_) {
return ((type == nvinfer1::DataType::kFLOAT || return ((type == nvinfer1::DataType::kFLOAT ||
type == nvinfer1::DataType::kHALF) && type == nvinfer1::DataType::kHALF) &&
...@@ -44,7 +44,7 @@ bool GeluPlugin::supportsFormat(nvinfer1::DataType type, ...@@ -44,7 +44,7 @@ bool GeluPlugin::supportsFormat(nvinfer1::DataType type,
nvinfer1::Dims GeluPlugin::getOutputDimensions(int index, nvinfer1::Dims GeluPlugin::getOutputDimensions(int index,
const nvinfer1::Dims* in_dims, const nvinfer1::Dims* in_dims,
int nb_inputs) { int nb_inputs) TRT_NOEXCEPT {
assert(nb_inputs == 1); assert(nb_inputs == 1);
assert(index < this->getNbOutputs()); assert(index < this->getNbOutputs());
nvinfer1::Dims const& input_dims = in_dims[0]; nvinfer1::Dims const& input_dims = in_dims[0];
...@@ -96,7 +96,8 @@ int GeluPlugin::enqueue(int batch_size, const void* const* inputs, ...@@ -96,7 +96,8 @@ int GeluPlugin::enqueue(int batch_size, const void* const* inputs,
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
void** outputs, void*, cudaStream_t stream) { void** outputs, void*, cudaStream_t stream) {
#else #else
void* const* outputs, void*, cudaStream_t stream) { void* const* outputs, void*,
cudaStream_t stream) TRT_NOEXCEPT {
#endif #endif
const auto& input_dims = this->getInputDims(0); const auto& input_dims = this->getInputDims(0);
int num = batch_size; int num = batch_size;
...@@ -132,13 +133,13 @@ int GeluPlugin::enqueue(int batch_size, const void* const* inputs, ...@@ -132,13 +133,13 @@ int GeluPlugin::enqueue(int batch_size, const void* const* inputs,
nvinfer1::DimsExprs GeluPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs GeluPluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) { nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT {
return inputs[0]; return inputs[0];
} }
bool GeluPluginDynamic::supportsFormatCombination( bool GeluPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc* in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc* in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of swish plugin shoule not be nullptr.")); "The input of swish plugin shoule not be nullptr."));
...@@ -167,7 +168,8 @@ bool GeluPluginDynamic::supportsFormatCombination( ...@@ -167,7 +168,8 @@ bool GeluPluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType GeluPluginDynamic::getOutputDataType( nvinfer1::DataType GeluPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType* input_types, int nb_inputs) const { int index, const nvinfer1::DataType* input_types,
int nb_inputs) const TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument( PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument(
"The Gelu Plugin only has one input, so the " "The Gelu Plugin only has one input, so the "
"index value should be 0, but get %d.", "index value should be 0, but get %d.",
...@@ -178,7 +180,8 @@ nvinfer1::DataType GeluPluginDynamic::getOutputDataType( ...@@ -178,7 +180,8 @@ nvinfer1::DataType GeluPluginDynamic::getOutputDataType(
int GeluPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc* input_desc, int GeluPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc* input_desc,
const nvinfer1::PluginTensorDesc* output_desc, const nvinfer1::PluginTensorDesc* output_desc,
const void* const* inputs, void* const* outputs, const void* const* inputs, void* const* outputs,
void* workspace, cudaStream_t stream) { void* workspace,
cudaStream_t stream) TRT_NOEXCEPT {
auto input_dims = input_desc[0].dims; auto input_dims = input_desc[0].dims;
size_t num = ProductDim(input_dims); size_t num = ProductDim(input_dims);
const int block_size = 256; const int block_size = 256;
......
paddle/fluid/inference/tensorrt/plugin/gelu_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -35,40 +35,48 @@ class GeluPlugin : public PluginTensorRT { ...@@ -35,40 +35,48 @@ class GeluPlugin : public PluginTensorRT {
} }
~GeluPlugin() {} ~GeluPlugin() {}
GeluPlugin* clone() const override { return new GeluPlugin(with_fp16_); } GeluPlugin* clone() const TRT_NOEXCEPT override {
return new GeluPlugin(with_fp16_);
}
const char* getPluginType() const override { return "gelu_plugin"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "gelu_plugin";
int initialize() override { return 0; } }
bool supportsFormat(nvinfer1::DataType type, int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
nvinfer1::PluginFormat format) const override; int initialize() TRT_NOEXCEPT override { return 0; }
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format)
const TRT_NOEXCEPT override;
nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs,
int nb_input_dims) override; int nb_input_dims) TRT_NOEXCEPT override;
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
int enqueue(int batch_size, const void* const* inputs, void** outputs, int enqueue(int batch_size, const void* const* inputs, void** outputs,
#else #else
int enqueue(int batch_size, const void* const* inputs, void* const* outputs, int enqueue(int batch_size, const void* const* inputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream) override; void* workspace, cudaStream_t stream) TRT_NOEXCEPT override;
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
return getBaseSerializationSize(); return getBaseSerializationSize();
} }
// TRT will call this func to serialize the configuration of TRT // TRT will call this func to serialize the configuration of TRT
// It should not be called by users. // It should not be called by users.
void serialize(void* buffer) const override { serializeBase(buffer); } void serialize(void* buffer) const TRT_NOEXCEPT override {
serializeBase(buffer);
}
}; };
class GeluPluginCreator : public TensorRTPluginCreator { class GeluPluginCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "gelu_plugin"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "gelu_plugin";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new GeluPlugin(serial_data, serial_length); return new GeluPlugin(serial_data, serial_length);
} }
}; };
...@@ -83,61 +91,66 @@ class GeluPluginDynamic : public DynamicPluginTensorRT { ...@@ -83,61 +91,66 @@ class GeluPluginDynamic : public DynamicPluginTensorRT {
} }
~GeluPluginDynamic() {} ~GeluPluginDynamic() {}
nvinfer1::IPluginV2DynamicExt* clone() const override { nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override {
return new GeluPluginDynamic(with_fp16_); return new GeluPluginDynamic(with_fp16_);
} }
const char* getPluginType() const override { return "gelu_plugin_dynamic"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "gelu_plugin_dynamic";
int initialize() override { return 0; } }
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
int initialize() TRT_NOEXCEPT override { return 0; }
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
return SerializedSize(with_fp16_); return SerializedSize(with_fp16_);
} }
void serialize(void* buffer) const override { void serialize(void* buffer) const TRT_NOEXCEPT override {
SerializeValue(&buffer, with_fp16_); SerializeValue(&buffer, with_fp16_);
} }
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) override; nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* in_out, const nvinfer1::PluginTensorDesc* in_out,
int nb_inputs, int nb_outputs) override; int nb_inputs,
int nb_outputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nb_inputs, int nb_inputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nb_outputs) override {} int nb_outputs) TRT_NOEXCEPT override {}
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nb_inputs, int nb_inputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nb_outputs) const override { int nb_outputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
int enqueue(const nvinfer1::PluginTensorDesc* input_desc, int enqueue(const nvinfer1::PluginTensorDesc* input_desc,
const nvinfer1::PluginTensorDesc* output_desc, const nvinfer1::PluginTensorDesc* output_desc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* input_types, int index, const nvinfer1::DataType* input_types,
int nb_inputs) const override; int nb_inputs) const TRT_NOEXCEPT override;
void destroy() override { delete this; } void destroy() TRT_NOEXCEPT override { delete this; }
}; };
class GeluPluginDynamicCreator : public TensorRTPluginCreator { class GeluPluginDynamicCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "gelu_plugin_dynamic"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "gelu_plugin_dynamic";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
auto plugin = new GeluPluginDynamic(serial_data, serial_length); auto plugin = new GeluPluginDynamic(serial_data, serial_length);
return plugin; return plugin;
} }
......
paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -22,7 +22,7 @@ namespace tensorrt { ...@@ -22,7 +22,7 @@ namespace tensorrt {
namespace plugin { namespace plugin {
nvinfer1::Dims HardSwishPlugin::getOutputDimensions( nvinfer1::Dims HardSwishPlugin::getOutputDimensions(
int index, const nvinfer1::Dims* in_dims, int nb_inputs) { int index, const nvinfer1::Dims* in_dims, int nb_inputs) TRT_NOEXCEPT {
assert(nb_inputs == 1); assert(nb_inputs == 1);
assert(index < this->getNbOutputs()); assert(index < this->getNbOutputs());
nvinfer1::Dims const& input_dims = in_dims[0]; nvinfer1::Dims const& input_dims = in_dims[0];
...@@ -54,7 +54,8 @@ int HardSwishPlugin::enqueue(int batch_size, const void* const* inputs, ...@@ -54,7 +54,8 @@ int HardSwishPlugin::enqueue(int batch_size, const void* const* inputs,
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
void** outputs, void*, cudaStream_t stream) { void** outputs, void*, cudaStream_t stream) {
#else #else
void* const* outputs, void*, cudaStream_t stream) { void* const* outputs, void*,
cudaStream_t stream) TRT_NOEXCEPT {
#endif #endif
const auto& input_dims = this->getInputDims(0); const auto& input_dims = this->getInputDims(0);
int num = batch_size; int num = batch_size;
......
paddle/fluid/inference/tensorrt/plugin/hard_swish_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -40,30 +40,32 @@ class HardSwishPlugin : public PluginTensorRT { ...@@ -40,30 +40,32 @@ class HardSwishPlugin : public PluginTensorRT {
} }
~HardSwishPlugin() {} ~HardSwishPlugin() {}
HardSwishPlugin* clone() const override { HardSwishPlugin* clone() const TRT_NOEXCEPT override {
return new HardSwishPlugin(threshold_, scale_, offset_); return new HardSwishPlugin(threshold_, scale_, offset_);
} }
const char* getPluginType() const override { return "hard_swish_plugin"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "hard_swish_plugin";
int initialize() override { return 0; } }
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
int initialize() TRT_NOEXCEPT override { return 0; }
nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs,
int nbInputDims) override; int nbInputDims) TRT_NOEXCEPT override;
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
int enqueue(int batchSize, const void* const* inputs, void** outputs, int enqueue(int batchSize, const void* const* inputs, void** outputs,
#else #else
int enqueue(int batchSize, const void* const* inputs, void* const* outputs, int enqueue(int batchSize, const void* const* inputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream) override; void* workspace, cudaStream_t stream) TRT_NOEXCEPT override;
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
return getBaseSerializationSize() + SerializedSize(threshold_) + return getBaseSerializationSize() + SerializedSize(threshold_) +
SerializedSize(scale_) + SerializedSize(offset_); SerializedSize(scale_) + SerializedSize(offset_);
} }
// TRT will call this func to serialize the configuration of TRT // TRT will call this func to serialize the configuration of TRT
// It should not be called by users. // It should not be called by users.
void serialize(void* buffer) const override { void serialize(void* buffer) const TRT_NOEXCEPT override {
serializeBase(buffer); serializeBase(buffer);
SerializeValue(&buffer, threshold_); SerializeValue(&buffer, threshold_);
SerializeValue(&buffer, scale_); SerializeValue(&buffer, scale_);
...@@ -78,13 +80,15 @@ class HardSwishPlugin : public PluginTensorRT { ...@@ -78,13 +80,15 @@ class HardSwishPlugin : public PluginTensorRT {
class HardSwishPluginCreator : public TensorRTPluginCreator { class HardSwishPluginCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "hard_swish_plugin"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "hard_swish_plugin";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new HardSwishPlugin(serial_data, serial_length); return new HardSwishPlugin(serial_data, serial_length);
} }
}; };
......
paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -39,10 +39,10 @@ cudnnStatus_t convert_trt2cudnn_dtype(nvinfer1::DataType trt_dtype, ...@@ -39,10 +39,10 @@ cudnnStatus_t convert_trt2cudnn_dtype(nvinfer1::DataType trt_dtype,
return CUDNN_STATUS_SUCCESS; return CUDNN_STATUS_SUCCESS;
} }
int InstanceNormPlugin::initialize() { return 0; } int InstanceNormPlugin::initialize() TRT_NOEXCEPT { return 0; }
nvinfer1::Dims InstanceNormPlugin::getOutputDimensions( nvinfer1::Dims InstanceNormPlugin::getOutputDimensions(
int index, const nvinfer1::Dims *inputDims, int nbInputs) { int index, const nvinfer1::Dims *inputDims, int nbInputs) TRT_NOEXCEPT {
assert(nbInputs == 1); assert(nbInputs == 1);
assert(index < this->getNbOutputs()); assert(index < this->getNbOutputs());
nvinfer1::Dims const &input_dims = inputDims[0]; nvinfer1::Dims const &input_dims = inputDims[0];
...@@ -50,8 +50,8 @@ nvinfer1::Dims InstanceNormPlugin::getOutputDimensions( ...@@ -50,8 +50,8 @@ nvinfer1::Dims InstanceNormPlugin::getOutputDimensions(
return output_dims; return output_dims;
} }
bool InstanceNormPlugin::supportsFormat(nvinfer1::DataType type, bool InstanceNormPlugin::supportsFormat(
nvinfer1::PluginFormat format) const { nvinfer1::DataType type, nvinfer1::PluginFormat format) const TRT_NOEXCEPT {
return ((type == nvinfer1::DataType::kFLOAT || return ((type == nvinfer1::DataType::kFLOAT ||
type == nvinfer1::DataType::kHALF) && type == nvinfer1::DataType::kHALF) &&
(format == nvinfer1::PluginFormat::kLINEAR)); (format == nvinfer1::PluginFormat::kLINEAR));
...@@ -63,7 +63,7 @@ int InstanceNormPlugin::enqueue(int batch_size, const void *const *inputs, ...@@ -63,7 +63,7 @@ int InstanceNormPlugin::enqueue(int batch_size, const void *const *inputs,
#else #else
void *const *outputs, void *workspace, void *const *outputs, void *workspace,
#endif #endif
cudaStream_t stream) { cudaStream_t stream) TRT_NOEXCEPT {
const auto &input_dims = this->getInputDims(0); const auto &input_dims = this->getInputDims(0);
PADDLE_ENFORCE_EQ(input_dims.nbDims, 3, PADDLE_ENFORCE_EQ(input_dims.nbDims, 3,
......
paddle/fluid/inference/tensorrt/plugin/instance_norm_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -39,7 +39,7 @@ class InstanceNormPlugin : public PluginTensorRT { ...@@ -39,7 +39,7 @@ class InstanceNormPlugin : public PluginTensorRT {
cudnnTensorDescriptor_t x_desc_, y_desc_, b_desc_; cudnnTensorDescriptor_t x_desc_, y_desc_, b_desc_;
public: public:
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
return getBaseSerializationSize() + SerializedSize(eps_) + return getBaseSerializationSize() + SerializedSize(eps_) +
SerializedSize(scale_) + SerializedSize(bias_); SerializedSize(scale_) + SerializedSize(bias_);
} }
...@@ -47,7 +47,7 @@ class InstanceNormPlugin : public PluginTensorRT { ...@@ -47,7 +47,7 @@ class InstanceNormPlugin : public PluginTensorRT {
// TRT will call this func when we need to serialize the configuration of // TRT will call this func when we need to serialize the configuration of
// tensorrt. // tensorrt.
// It should not be called by users. // It should not be called by users.
void serialize(void *buffer) const override { void serialize(void *buffer) const TRT_NOEXCEPT override {
serializeBase(buffer); serializeBase(buffer);
SerializeValue(&buffer, eps_); SerializeValue(&buffer, eps_);
SerializeValue(&buffer, scale_); SerializeValue(&buffer, scale_);
...@@ -89,37 +89,41 @@ class InstanceNormPlugin : public PluginTensorRT { ...@@ -89,37 +89,41 @@ class InstanceNormPlugin : public PluginTensorRT {
platform::dynload::cudnnDestroyTensorDescriptor(b_desc_); platform::dynload::cudnnDestroyTensorDescriptor(b_desc_);
} }
int initialize() override; int initialize() TRT_NOEXCEPT override;
InstanceNormPlugin *clone() const override { InstanceNormPlugin *clone() const TRT_NOEXCEPT override {
return new InstanceNormPlugin(eps_, scale_, bias_); return new InstanceNormPlugin(eps_, scale_, bias_);
} }
const char *getPluginType() const override { return "instance_norm_plugin"; } const char *getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "instance_norm_plugin";
}
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims *inputs, nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims *inputs,
int nbInputDims) override; int nbInputDims) TRT_NOEXCEPT override;
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
int enqueue(int batchSize, const void *const *inputs, void **outputs, int enqueue(int batchSize, const void *const *inputs, void **outputs,
#else #else
int enqueue(int batchSize, const void *const *inputs, void *const *outputs, int enqueue(int batchSize, const void *const *inputs, void *const *outputs,
#endif #endif
void *workspace, cudaStream_t stream) override; void *workspace, cudaStream_t stream) TRT_NOEXCEPT override;
bool supportsFormat(nvinfer1::DataType type, bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format)
nvinfer1::PluginFormat format) const override; const TRT_NOEXCEPT override;
}; };
class InstanceNormPluginCreator : public TensorRTPluginCreator { class InstanceNormPluginCreator : public TensorRTPluginCreator {
public: public:
const char *getPluginName() const override { return "instance_norm_plugin"; } const char *getPluginName() const TRT_NOEXCEPT override {
return "instance_norm_plugin";
}
const char *getPluginVersion() const override { return "1"; } const char *getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2 *deserializePlugin(const char *name, nvinfer1::IPluginV2 *deserializePlugin(
const void *serial_data, const char *name, const void *serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new InstanceNormPlugin(serial_data, serial_length); return new InstanceNormPlugin(serial_data, serial_length);
} }
}; };
......
paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -24,10 +24,10 @@ namespace inference { ...@@ -24,10 +24,10 @@ namespace inference {
namespace tensorrt { namespace tensorrt {
namespace plugin { namespace plugin {
int LayerNormPlugin::initialize() { return 0; } int LayerNormPlugin::initialize() TRT_NOEXCEPT { return 0; }
nvinfer1::Dims LayerNormPlugin::getOutputDimensions( nvinfer1::Dims LayerNormPlugin::getOutputDimensions(
int index, const nvinfer1::Dims *inputDims, int nbInputs) { int index, const nvinfer1::Dims *inputDims, int nbInputs) TRT_NOEXCEPT {
assert(nbInputs == 1); assert(nbInputs == 1);
assert(index < this->getNbOutputs()); assert(index < this->getNbOutputs());
nvinfer1::Dims const &input_dims = inputDims[0]; nvinfer1::Dims const &input_dims = inputDims[0];
...@@ -41,10 +41,10 @@ int LayerNormPlugin::enqueue(int batch_size, const void *const *inputs, ...@@ -41,10 +41,10 @@ int LayerNormPlugin::enqueue(int batch_size, const void *const *inputs,
#else #else
void *const *outputs, void *workspace, void *const *outputs, void *workspace,
#endif #endif
cudaStream_t stream) { cudaStream_t stream) TRT_NOEXCEPT {
const auto &input_dims = this->getInputDims(0); const auto &input_dims = this->getInputDims(0);
const float *input = reinterpret_cast<const float *>(inputs[0]); const float *input = reinterpret_cast<const float *>(inputs[0]);
float *output = reinterpret_cast<float **>(outputs)[0]; float *output = reinterpret_cast<float *const *>(outputs)[0];
int begin_norm_axis = begin_norm_axis_; int begin_norm_axis = begin_norm_axis_;
float eps = eps_; float eps = eps_;
...@@ -91,13 +91,13 @@ int LayerNormPlugin::enqueue(int batch_size, const void *const *inputs, ...@@ -91,13 +91,13 @@ int LayerNormPlugin::enqueue(int batch_size, const void *const *inputs,
nvinfer1::DimsExprs LayerNormPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs LayerNormPluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs *inputDims, int nb_inputs, int output_index, const nvinfer1::DimsExprs *inputDims, int nb_inputs,
nvinfer1::IExprBuilder &expr_builder) { nvinfer1::IExprBuilder &expr_builder) TRT_NOEXCEPT {
return inputDims[0]; return inputDims[0];
} }
bool LayerNormPluginDynamic::supportsFormatCombination( bool LayerNormPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of layernorm plugin shoule not be nullptr.")); "The input of layernorm plugin shoule not be nullptr."));
...@@ -118,7 +118,8 @@ bool LayerNormPluginDynamic::supportsFormatCombination( ...@@ -118,7 +118,8 @@ bool LayerNormPluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType LayerNormPluginDynamic::getOutputDataType( nvinfer1::DataType LayerNormPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType *input_types, int nb_inputs) const { int index, const nvinfer1::DataType *input_types,
int nb_inputs) const TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(index, 0, PADDLE_ENFORCE_EQ(index, 0,
platform::errors::InvalidArgument( platform::errors::InvalidArgument(
"The LayerNormPlugin only has one input, so the " "The LayerNormPlugin only has one input, so the "
...@@ -130,7 +131,7 @@ nvinfer1::DataType LayerNormPluginDynamic::getOutputDataType( ...@@ -130,7 +131,7 @@ nvinfer1::DataType LayerNormPluginDynamic::getOutputDataType(
int LayerNormPluginDynamic::enqueue( int LayerNormPluginDynamic::enqueue(
const nvinfer1::PluginTensorDesc *input_desc, const nvinfer1::PluginTensorDesc *input_desc,
const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs, const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs,
void *const *outputs, void *workspace, cudaStream_t stream) { void *const *outputs, void *workspace, cudaStream_t stream) TRT_NOEXCEPT {
const auto &input_dims = input_desc[0].dims; const auto &input_dims = input_desc[0].dims;
int begin_norm_axis = begin_norm_axis_; int begin_norm_axis = begin_norm_axis_;
float eps = eps_; float eps = eps_;
......
paddle/fluid/inference/tensorrt/plugin/layer_norm_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -40,7 +40,7 @@ class LayerNormPlugin : public PluginTensorRT { ...@@ -40,7 +40,7 @@ class LayerNormPlugin : public PluginTensorRT {
std::vector<int64_t> variance_shape_; std::vector<int64_t> variance_shape_;
public: public:
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
return getBaseSerializationSize() + SerializedSize(bias_) + return getBaseSerializationSize() + SerializedSize(bias_) +
SerializedSize(scale_) + SerializedSize(begin_norm_axis_) + SerializedSize(scale_) + SerializedSize(begin_norm_axis_) +
SerializedSize(eps_) + SerializedSize(mean_shape_) + SerializedSize(eps_) + SerializedSize(mean_shape_) +
...@@ -50,7 +50,7 @@ class LayerNormPlugin : public PluginTensorRT { ...@@ -50,7 +50,7 @@ class LayerNormPlugin : public PluginTensorRT {
// TRT will call this func when we need to serialize the configuration of // TRT will call this func when we need to serialize the configuration of
// tensorrt. // tensorrt.
// It should not be called by users. // It should not be called by users.
void serialize(void* buffer) const override { void serialize(void* buffer) const TRT_NOEXCEPT override {
serializeBase(buffer); serializeBase(buffer);
SerializeValue(&buffer, bias_); SerializeValue(&buffer, bias_);
SerializeValue(&buffer, scale_); SerializeValue(&buffer, scale_);
...@@ -86,35 +86,39 @@ class LayerNormPlugin : public PluginTensorRT { ...@@ -86,35 +86,39 @@ class LayerNormPlugin : public PluginTensorRT {
DeserializeValue(&serialData, &serialLength, &variance_shape_); DeserializeValue(&serialData, &serialLength, &variance_shape_);
} }
~LayerNormPlugin() {} ~LayerNormPlugin() {}
int initialize() override; int initialize() TRT_NOEXCEPT override;
LayerNormPlugin* clone() const override { LayerNormPlugin* clone() const TRT_NOEXCEPT override {
return new LayerNormPlugin(bias_.data(), bias_.size(), scale_.data(), return new LayerNormPlugin(bias_.data(), bias_.size(), scale_.data(),
scale_.size(), begin_norm_axis_, eps_, scale_.size(), begin_norm_axis_, eps_,
mean_shape_, variance_shape_); mean_shape_, variance_shape_);
} }
const char* getPluginType() const override { return "layernorm_plugin"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "layernorm_plugin";
}
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs,
int nbInputDims) override; int nbInputDims) TRT_NOEXCEPT override;
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
int enqueue(int batchSize, const void* const* inputs, void** outputs, int enqueue(int batchSize, const void* const* inputs, void** outputs,
#else #else
int enqueue(int batchSize, const void* const* inputs, void* const* outputs, int enqueue(int batchSize, const void* const* inputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream) override; void* workspace, cudaStream_t stream) TRT_NOEXCEPT override;
}; };
class LayerNormPluginCreator : public TensorRTPluginCreator { class LayerNormPluginCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "layernorm_plugin"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "layernorm_plugin";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new LayerNormPlugin(serial_data, serial_length); return new LayerNormPlugin(serial_data, serial_length);
} }
}; };
...@@ -145,25 +149,25 @@ class LayerNormPluginDynamic : public DynamicPluginTensorRT { ...@@ -145,25 +149,25 @@ class LayerNormPluginDynamic : public DynamicPluginTensorRT {
DeserializeValue(&serialData, &serialLength, &mean_shape_); DeserializeValue(&serialData, &serialLength, &mean_shape_);
DeserializeValue(&serialData, &serialLength, &variance_shape_); DeserializeValue(&serialData, &serialLength, &variance_shape_);
} }
nvinfer1::IPluginV2DynamicExt* clone() const override { nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override {
return new LayerNormPluginDynamic(bias_.data(), bias_.size(), scale_.data(), return new LayerNormPluginDynamic(bias_.data(), bias_.size(), scale_.data(),
scale_.size(), begin_norm_axis_, eps_, scale_.size(), begin_norm_axis_, eps_,
mean_shape_, variance_shape_); mean_shape_, variance_shape_);
} }
const char* getPluginType() const override { const char* getPluginType() const TRT_NOEXCEPT override {
return "layernorm_plugin_dynamic"; return "layernorm_plugin_dynamic";
} }
int getNbOutputs() const override { return 1; } int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
int initialize() override { return 0; } int initialize() TRT_NOEXCEPT override { return 0; }
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
return SerializedSize(bias_) + SerializedSize(scale_) + return SerializedSize(bias_) + SerializedSize(scale_) +
SerializedSize(begin_norm_axis_) + SerializedSize(eps_) + SerializedSize(begin_norm_axis_) + SerializedSize(eps_) +
SerializedSize(mean_shape_) + SerializedSize(variance_shape_); SerializedSize(mean_shape_) + SerializedSize(variance_shape_);
} }
void serialize(void* buffer) const override { void serialize(void* buffer) const TRT_NOEXCEPT override {
SerializeValue(&buffer, bias_); SerializeValue(&buffer, bias_);
SerializeValue(&buffer, scale_); SerializeValue(&buffer, scale_);
SerializeValue(&buffer, begin_norm_axis_); SerializeValue(&buffer, begin_norm_axis_);
...@@ -174,33 +178,34 @@ class LayerNormPluginDynamic : public DynamicPluginTensorRT { ...@@ -174,33 +178,34 @@ class LayerNormPluginDynamic : public DynamicPluginTensorRT {
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) override; nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* inOut, const nvinfer1::PluginTensorDesc* inOut,
int nbInputs, int nbOutputs) override; int nbInputs,
int nbOutputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nbInputs, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nbOutputs) override {} int nbOutputs) TRT_NOEXCEPT override {}
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nbInputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const override { int nbOutputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* inputTypes, int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const override; int nbInputs) const TRT_NOEXCEPT override;
void destroy() override { delete this; } void destroy() TRT_NOEXCEPT override { delete this; }
private: private:
std::vector<float> bias_; std::vector<float> bias_;
...@@ -217,15 +222,15 @@ class LayerNormPluginDynamic : public DynamicPluginTensorRT { ...@@ -217,15 +222,15 @@ class LayerNormPluginDynamic : public DynamicPluginTensorRT {
class LayerNormPluginDynamicCreator : public TensorRTPluginCreator { class LayerNormPluginDynamicCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { const char* getPluginName() const TRT_NOEXCEPT override {
return "layernorm_plugin_dynamic"; return "layernorm_plugin_dynamic";
} }
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new LayerNormPluginDynamic(serial_data, serial_length); return new LayerNormPluginDynamic(serial_data, serial_length);
} }
}; };
......
paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -22,7 +22,7 @@ namespace plugin { ...@@ -22,7 +22,7 @@ namespace plugin {
nvinfer1::Dims PoolPlugin::getOutputDimensions(int index, nvinfer1::Dims PoolPlugin::getOutputDimensions(int index,
const nvinfer1::Dims *inputDims, const nvinfer1::Dims *inputDims,
int nbInputs) { int nbInputs) TRT_NOEXCEPT {
assert(nbInputs == 1); assert(nbInputs == 1);
assert(index == 0); assert(index == 0);
assert(inputDims[0].nbDims == 3); assert(inputDims[0].nbDims == 3);
...@@ -37,15 +37,16 @@ nvinfer1::Dims PoolPlugin::getOutputDimensions(int index, ...@@ -37,15 +37,16 @@ nvinfer1::Dims PoolPlugin::getOutputDimensions(int index,
int PoolPlugin::enqueue(int batchSize, const void *const *inputs, int PoolPlugin::enqueue(int batchSize, const void *const *inputs,
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
void **outputs, void *workspace, cudaStream_t stream) { void **outputs, void *workspace,
cudaStream_t stream) TRT_NOEXCEPT {
#else #else
void *const *outputs, void *workspace, void *const *outputs, void *workspace,
cudaStream_t stream) { cudaStream_t stream) TRT_NOEXCEPT {
#endif #endif
auto const &input_dims = this->getInputDims(0); auto const &input_dims = this->getInputDims(0);
int input_size = 0; int input_size = 0;
float const *idata = reinterpret_cast<float const *>(inputs[0]); float const *idata = reinterpret_cast<float const *>(inputs[0]);
float **odatas = reinterpret_cast<float **>(outputs); float *const *odatas = reinterpret_cast<float *const *>(outputs);
std::vector<int> input_shape = input_shape_; std::vector<int> input_shape = input_shape_;
std::vector<int> output_shape = output_shape_; std::vector<int> output_shape = output_shape_;
...@@ -87,14 +88,14 @@ PoolPluginDynamic::PoolPluginDynamic(void const *serialData, ...@@ -87,14 +88,14 @@ PoolPluginDynamic::PoolPluginDynamic(void const *serialData,
DeserializeValue(&serialData, &serialLength, &is_global_); DeserializeValue(&serialData, &serialLength, &is_global_);
} }
size_t PoolPluginDynamic::getSerializationSize() const { size_t PoolPluginDynamic::getSerializationSize() const TRT_NOEXCEPT {
return SerializedSize(ceil_mode_) + SerializedSize(pool_type_.c_str()) + return SerializedSize(ceil_mode_) + SerializedSize(pool_type_.c_str()) +
SerializedSize(adaptive_) + SerializedSize(ksize_) + SerializedSize(adaptive_) + SerializedSize(ksize_) +
SerializedSize(strides_) + SerializedSize(paddings_) + SerializedSize(strides_) + SerializedSize(paddings_) +
SerializedSize(is_global_); SerializedSize(is_global_);
} }
void PoolPluginDynamic::serialize(void *buffer) const { void PoolPluginDynamic::serialize(void *buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, ceil_mode_); SerializeValue(&buffer, ceil_mode_);
SerializeValue(&buffer, pool_type_.c_str()); SerializeValue(&buffer, pool_type_.c_str());
SerializeValue(&buffer, adaptive_); SerializeValue(&buffer, adaptive_);
...@@ -106,7 +107,7 @@ void PoolPluginDynamic::serialize(void *buffer) const { ...@@ -106,7 +107,7 @@ void PoolPluginDynamic::serialize(void *buffer) const {
nvinfer1::DimsExprs PoolPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs PoolPluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs,
nvinfer1::IExprBuilder &expr_builder) { nvinfer1::IExprBuilder &expr_builder) TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(nb_inputs, 1, PADDLE_ENFORCE_EQ(nb_inputs, 1,
platform::errors::InvalidArgument( platform::errors::InvalidArgument(
"The Split plugin should be only one input.")); "The Split plugin should be only one input."));
...@@ -181,7 +182,7 @@ nvinfer1::DimsExprs PoolPluginDynamic::getOutputDimensions( ...@@ -181,7 +182,7 @@ nvinfer1::DimsExprs PoolPluginDynamic::getOutputDimensions(
bool PoolPluginDynamic::supportsFormatCombination( bool PoolPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of swish plugin shoule not be nullptr.")); "The input of swish plugin shoule not be nullptr."));
...@@ -198,7 +199,8 @@ bool PoolPluginDynamic::supportsFormatCombination( ...@@ -198,7 +199,8 @@ bool PoolPluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType PoolPluginDynamic::getOutputDataType( nvinfer1::DataType PoolPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType *input_types, int nb_inputs) const { int index, const nvinfer1::DataType *input_types,
int nb_inputs) const TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument( PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument(
"The Pool Plugin only has one input, so the " "The Pool Plugin only has one input, so the "
"index value should be 0, but get %d.", "index value should be 0, but get %d.",
...@@ -212,7 +214,8 @@ nvinfer1::DataType PoolPluginDynamic::getOutputDataType( ...@@ -212,7 +214,8 @@ nvinfer1::DataType PoolPluginDynamic::getOutputDataType(
int PoolPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc *input_desc, int PoolPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc *input_desc,
const nvinfer1::PluginTensorDesc *output_desc, const nvinfer1::PluginTensorDesc *output_desc,
const void *const *inputs, void *const *outputs, const void *const *inputs, void *const *outputs,
void *workspace, cudaStream_t stream) { void *workspace,
cudaStream_t stream) TRT_NOEXCEPT {
auto input_dims = input_desc[0].dims; auto input_dims = input_desc[0].dims;
int n = input_dims.d[0]; int n = input_dims.d[0];
int c = input_dims.d[1]; int c = input_dims.d[1];
......
paddle/fluid/inference/tensorrt/plugin/pool_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -57,7 +57,7 @@ static std::vector<int> CalcOutputSize(const std::vector<int>& input_shape, ...@@ -57,7 +57,7 @@ static std::vector<int> CalcOutputSize(const std::vector<int>& input_shape,
class PoolPlugin : public PluginTensorRT { class PoolPlugin : public PluginTensorRT {
public: public:
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
return getBaseSerializationSize() + SerializedSize(ceil_mode_) + return getBaseSerializationSize() + SerializedSize(ceil_mode_) +
SerializedSize(pool_type_) + SerializedSize(adaptive_) + SerializedSize(pool_type_) + SerializedSize(adaptive_) +
SerializedSize(ksize_) + SerializedSize(strides_) + SerializedSize(ksize_) + SerializedSize(strides_) +
...@@ -67,7 +67,7 @@ class PoolPlugin : public PluginTensorRT { ...@@ -67,7 +67,7 @@ class PoolPlugin : public PluginTensorRT {
// TRT will call this func when we need to serialize the configuration of // TRT will call this func when we need to serialize the configuration of
// tensorrt. // tensorrt.
void serialize(void* buffer) const override { void serialize(void* buffer) const TRT_NOEXCEPT override {
serializeBase(buffer); serializeBase(buffer);
SerializeValue(&buffer, ceil_mode_); SerializeValue(&buffer, ceil_mode_);
SerializeValue(&buffer, pool_type_); SerializeValue(&buffer, pool_type_);
...@@ -116,22 +116,24 @@ class PoolPlugin : public PluginTensorRT { ...@@ -116,22 +116,24 @@ class PoolPlugin : public PluginTensorRT {
DeserializeValue(&serialData, &serialLength, &output_shape_); DeserializeValue(&serialData, &serialLength, &output_shape_);
} }
PoolPlugin* clone() const override { PoolPlugin* clone() const TRT_NOEXCEPT override {
return new PoolPlugin(ceil_mode_, pool_type_, adaptive_, ksize_, strides_, return new PoolPlugin(ceil_mode_, pool_type_, adaptive_, ksize_, strides_,
paddings_, input_shape_); paddings_, input_shape_);
} }
const char* getPluginType() const override { return "pool_plugin"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "pool_plugin";
}
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs,
int nbInputDims) override; int nbInputDims) TRT_NOEXCEPT override;
int initialize() override { return 0; } int initialize() TRT_NOEXCEPT override { return 0; }
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
int enqueue(int batchSize, const void* const* inputs, void** outputs, int enqueue(int batchSize, const void* const* inputs, void** outputs,
#else #else
int enqueue(int batchSize, const void* const* inputs, void* const* outputs, int enqueue(int batchSize, const void* const* inputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream) override; void* workspace, cudaStream_t stream) TRT_NOEXCEPT override;
private: private:
bool ceil_mode_; bool ceil_mode_;
...@@ -146,13 +148,15 @@ class PoolPlugin : public PluginTensorRT { ...@@ -146,13 +148,15 @@ class PoolPlugin : public PluginTensorRT {
class PoolPluginCreator : public TensorRTPluginCreator { class PoolPluginCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "pool_plugin"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "pool_plugin";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new PoolPlugin(serial_data, serial_length); return new PoolPlugin(serial_data, serial_length);
} }
}; };
...@@ -176,47 +180,50 @@ class PoolPluginDynamic : public DynamicPluginTensorRT { ...@@ -176,47 +180,50 @@ class PoolPluginDynamic : public DynamicPluginTensorRT {
PoolPluginDynamic(void const* serialData, size_t serialLength); PoolPluginDynamic(void const* serialData, size_t serialLength);
~PoolPluginDynamic() {} ~PoolPluginDynamic() {}
nvinfer1::IPluginV2DynamicExt* clone() const override { nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override {
return new PoolPluginDynamic(ceil_mode_, pool_type_, adaptive_, ksize_, return new PoolPluginDynamic(ceil_mode_, pool_type_, adaptive_, ksize_,
strides_, paddings_, is_global_); strides_, paddings_, is_global_);
} }
const char* getPluginType() const override { return "pool_plugin_dynamic"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "pool_plugin_dynamic";
int initialize() override { return 0; } }
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
int initialize() TRT_NOEXCEPT override { return 0; }
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) override; nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* inOut, const nvinfer1::PluginTensorDesc* inOut,
int nbInputs, int nbOutputs) override; int nbInputs,
int nbOutputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nbInputs, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nbOutputs) override {} int nbOutputs) TRT_NOEXCEPT override {}
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nbInputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const override { int nbOutputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* inputTypes, int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const override; int nbInputs) const TRT_NOEXCEPT override;
void destroy() override { delete this; } void destroy() TRT_NOEXCEPT override { delete this; }
private: private:
bool ceil_mode_; bool ceil_mode_;
...@@ -230,13 +237,15 @@ class PoolPluginDynamic : public DynamicPluginTensorRT { ...@@ -230,13 +237,15 @@ class PoolPluginDynamic : public DynamicPluginTensorRT {
class PoolPluginDynamicCreator : public TensorRTPluginCreator { class PoolPluginDynamicCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "pool_plugin_dynamic"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "pool_plugin_dynamic";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new PoolPluginDynamic(serial_data, serial_length); return new PoolPluginDynamic(serial_data, serial_length);
} }
}; };
......
paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -26,14 +26,14 @@ namespace inference { ...@@ -26,14 +26,14 @@ namespace inference {
namespace tensorrt { namespace tensorrt {
namespace plugin { namespace plugin {
int PReluPlugin::initialize() { int PReluPlugin::initialize() TRT_NOEXCEPT {
cudaMalloc(&p_gpu_weight_, sizeof(float) * weight_.size()); cudaMalloc(&p_gpu_weight_, sizeof(float) * weight_.size());
cudaMemcpy(p_gpu_weight_, weight_.data(), weight_.size() * sizeof(float), cudaMemcpy(p_gpu_weight_, weight_.data(), weight_.size() * sizeof(float),
cudaMemcpyHostToDevice); cudaMemcpyHostToDevice);
return 0; return 0;
} }
void PReluPlugin::terminate() { void PReluPlugin::terminate() TRT_NOEXCEPT {
if (p_gpu_weight_) { if (p_gpu_weight_) {
cudaFree(p_gpu_weight_); cudaFree(p_gpu_weight_);
p_gpu_weight_ = nullptr; p_gpu_weight_ = nullptr;
...@@ -42,7 +42,7 @@ void PReluPlugin::terminate() { ...@@ -42,7 +42,7 @@ void PReluPlugin::terminate() {
nvinfer1::Dims PReluPlugin::getOutputDimensions(int index, nvinfer1::Dims PReluPlugin::getOutputDimensions(int index,
const nvinfer1::Dims *inputDims, const nvinfer1::Dims *inputDims,
int nbInputs) { int nbInputs) TRT_NOEXCEPT {
assert(nbInputs == 1); assert(nbInputs == 1);
assert(index < this->getNbOutputs()); assert(index < this->getNbOutputs());
nvinfer1::Dims const &input_dims = inputDims[0]; nvinfer1::Dims const &input_dims = inputDims[0];
...@@ -55,14 +55,14 @@ int PReluPlugin::enqueue(int batch_size, const void *const *inputs, ...@@ -55,14 +55,14 @@ int PReluPlugin::enqueue(int batch_size, const void *const *inputs,
void **outputs, void *workspace, cudaStream_t stream) { void **outputs, void *workspace, cudaStream_t stream) {
#else #else
void *const *outputs, void *workspace, void *const *outputs, void *workspace,
cudaStream_t stream) { cudaStream_t stream) TRT_NOEXCEPT {
#endif #endif
// input dims is CHW. // input dims is CHW.
const auto &input_dims = this->getInputDims(0); const auto &input_dims = this->getInputDims(0);
const float *input = reinterpret_cast<const float *>(inputs[0]); const float *input = reinterpret_cast<const float *>(inputs[0]);
// const float *alpha = reinterpret_cast<const float *>(alpha_.get().values); // const float *alpha = reinterpret_cast<const float *>(alpha_.get().values);
const float *alpha = p_gpu_weight_; const float *alpha = p_gpu_weight_;
float *output = reinterpret_cast<float **>(outputs)[0]; float *const output = reinterpret_cast<float *const *>(outputs)[0];
int numel = 1; int numel = 1;
for (int i = 0; i < input_dims.nbDims; i++) { for (int i = 0; i < input_dims.nbDims; i++) {
numel *= input_dims.d[i]; numel *= input_dims.d[i];
...@@ -86,13 +86,13 @@ int PReluPlugin::enqueue(int batch_size, const void *const *inputs, ...@@ -86,13 +86,13 @@ int PReluPlugin::enqueue(int batch_size, const void *const *inputs,
#if IS_TRT_VERSION_GE(6000) #if IS_TRT_VERSION_GE(6000)
void PReluPluginDynamic::terminate() { void PReluPluginDynamic::terminate() TRT_NOEXCEPT {
if (p_gpu_weight_) { if (p_gpu_weight_) {
cudaFree(p_gpu_weight_); cudaFree(p_gpu_weight_);
} }
} }
int PReluPluginDynamic::initialize() { int PReluPluginDynamic::initialize() TRT_NOEXCEPT {
cudaMalloc(&p_gpu_weight_, sizeof(float) * weight_.size()); cudaMalloc(&p_gpu_weight_, sizeof(float) * weight_.size());
cudaMemcpy(p_gpu_weight_, weight_.data(), weight_.size() * sizeof(float), cudaMemcpy(p_gpu_weight_, weight_.data(), weight_.size() * sizeof(float),
cudaMemcpyHostToDevice); cudaMemcpyHostToDevice);
...@@ -107,24 +107,24 @@ PReluPluginDynamic::PReluPluginDynamic(void const *serialData, ...@@ -107,24 +107,24 @@ PReluPluginDynamic::PReluPluginDynamic(void const *serialData,
mode_ = std::string(prelu_mode); mode_ = std::string(prelu_mode);
} }
size_t PReluPluginDynamic::getSerializationSize() const { size_t PReluPluginDynamic::getSerializationSize() const TRT_NOEXCEPT {
return SerializedSize(mode_.c_str()) + SerializedSize(weight_); return SerializedSize(mode_.c_str()) + SerializedSize(weight_);
} }
void PReluPluginDynamic::serialize(void *buffer) const { void PReluPluginDynamic::serialize(void *buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, weight_); SerializeValue(&buffer, weight_);
SerializeValue(&buffer, mode_.c_str()); SerializeValue(&buffer, mode_.c_str());
} }
nvinfer1::DimsExprs PReluPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs PReluPluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs,
nvinfer1::IExprBuilder &expr_builder) { nvinfer1::IExprBuilder &expr_builder) TRT_NOEXCEPT {
return inputs[0]; return inputs[0];
} }
bool PReluPluginDynamic::supportsFormatCombination( bool PReluPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of swish plugin shoule not be nullptr.")); "The input of swish plugin shoule not be nullptr."));
...@@ -141,7 +141,8 @@ bool PReluPluginDynamic::supportsFormatCombination( ...@@ -141,7 +141,8 @@ bool PReluPluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType PReluPluginDynamic::getOutputDataType( nvinfer1::DataType PReluPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType *input_types, int nb_inputs) const { int index, const nvinfer1::DataType *input_types,
int nb_inputs) const TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument( PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument(
"The PRelu Plugin only has one input, so the " "The PRelu Plugin only has one input, so the "
"index value should be 0, but get %d.", "index value should be 0, but get %d.",
...@@ -155,7 +156,8 @@ nvinfer1::DataType PReluPluginDynamic::getOutputDataType( ...@@ -155,7 +156,8 @@ nvinfer1::DataType PReluPluginDynamic::getOutputDataType(
int PReluPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc *input_desc, int PReluPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc *input_desc,
const nvinfer1::PluginTensorDesc *output_desc, const nvinfer1::PluginTensorDesc *output_desc,
const void *const *inputs, void *const *outputs, const void *const *inputs, void *const *outputs,
void *workspace, cudaStream_t stream) { void *workspace,
cudaStream_t stream) TRT_NOEXCEPT {
auto input_dims = input_desc[0].dims; auto input_dims = input_desc[0].dims;
const float *alpha = p_gpu_weight_; const float *alpha = p_gpu_weight_;
const float *input = static_cast<const float *>(inputs[0]); const float *input = static_cast<const float *>(inputs[0]);
......
paddle/fluid/inference/tensorrt/plugin/prelu_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -34,7 +34,7 @@ class PReluPlugin : public PluginTensorRT { ...@@ -34,7 +34,7 @@ class PReluPlugin : public PluginTensorRT {
std::string mode_; std::string mode_;
public: public:
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
return getBaseSerializationSize() + SerializedSize(mode_.c_str()) + return getBaseSerializationSize() + SerializedSize(mode_.c_str()) +
SerializedSize(weight_); SerializedSize(weight_);
} }
...@@ -42,7 +42,7 @@ class PReluPlugin : public PluginTensorRT { ...@@ -42,7 +42,7 @@ class PReluPlugin : public PluginTensorRT {
// TRT will call this func when we need to serialize the configuration of // TRT will call this func when we need to serialize the configuration of
// tensorrt. // tensorrt.
// It should not be called by users. // It should not be called by users.
void serialize(void* buffer) const override { void serialize(void* buffer) const TRT_NOEXCEPT override {
serializeBase(buffer); serializeBase(buffer);
SerializeValue(&buffer, weight_); SerializeValue(&buffer, weight_);
SerializeValue(&buffer, mode_.c_str()); SerializeValue(&buffer, mode_.c_str());
...@@ -65,36 +65,40 @@ class PReluPlugin : public PluginTensorRT { ...@@ -65,36 +65,40 @@ class PReluPlugin : public PluginTensorRT {
mode_ = std::string(prelu_mode); mode_ = std::string(prelu_mode);
} }
~PReluPlugin() {} ~PReluPlugin() {}
int initialize() override; int initialize() TRT_NOEXCEPT override;
void terminate() override; void terminate() TRT_NOEXCEPT override;
PReluPlugin* clone() const override { PReluPlugin* clone() const TRT_NOEXCEPT override {
auto* ptr = new PReluPlugin(weight_.data(), weight_.size(), mode_); auto* ptr = new PReluPlugin(weight_.data(), weight_.size(), mode_);
ptr->p_gpu_weight_ = p_gpu_weight_; ptr->p_gpu_weight_ = p_gpu_weight_;
return ptr; return ptr;
} }
const char* getPluginType() const override { return "prelu_plugin"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "prelu_plugin";
}
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs,
int nbInputDims) override; int nbInputDims) TRT_NOEXCEPT override;
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
int enqueue(int batchSize, const void* const* inputs, void** outputs, int enqueue(int batchSize, const void* const* inputs, void** outputs,
#else #else
int enqueue(int batchSize, const void* const* inputs, void* const* outputs, int enqueue(int batchSize, const void* const* inputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream) override; void* workspace, cudaStream_t stream) TRT_NOEXCEPT override;
}; };
class PReluPluginCreator : public TensorRTPluginCreator { class PReluPluginCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "prelu_plugin"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "prelu_plugin";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new PReluPlugin(serial_data, serial_length); return new PReluPlugin(serial_data, serial_length);
} }
}; };
...@@ -112,49 +116,52 @@ class PReluPluginDynamic : public DynamicPluginTensorRT { ...@@ -112,49 +116,52 @@ class PReluPluginDynamic : public DynamicPluginTensorRT {
PReluPluginDynamic(void const* serialData, size_t serialLength); PReluPluginDynamic(void const* serialData, size_t serialLength);
~PReluPluginDynamic() {} ~PReluPluginDynamic() {}
nvinfer1::IPluginV2DynamicExt* clone() const override { nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override {
auto ptr = new PReluPluginDynamic(weight_.data(), weight_.size(), mode_); auto ptr = new PReluPluginDynamic(weight_.data(), weight_.size(), mode_);
ptr->p_gpu_weight_ = p_gpu_weight_; ptr->p_gpu_weight_ = p_gpu_weight_;
return ptr; return ptr;
} }
const char* getPluginType() const override { return "prelu_plugin_dynamic"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "prelu_plugin_dynamic";
int initialize() override; }
void terminate() override; int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
int initialize() TRT_NOEXCEPT override;
void terminate() TRT_NOEXCEPT override;
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) override; nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* inOut, const nvinfer1::PluginTensorDesc* inOut,
int nbInputs, int nbOutputs) override; int nbInputs,
int nbOutputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nbInputs, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nbOutputs) override {} int nbOutputs) TRT_NOEXCEPT override {}
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nbInputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const override { int nbOutputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* inputTypes, int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const override; int nbInputs) const TRT_NOEXCEPT override;
void destroy() override { delete this; } void destroy() TRT_NOEXCEPT override { delete this; }
private: private:
std::vector<float> weight_; std::vector<float> weight_;
...@@ -165,13 +172,15 @@ class PReluPluginDynamic : public DynamicPluginTensorRT { ...@@ -165,13 +172,15 @@ class PReluPluginDynamic : public DynamicPluginTensorRT {
class PReluPluginDynamicCreator : public TensorRTPluginCreator { class PReluPluginDynamicCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "prelu_plugin_dynamic"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "prelu_plugin_dynamic";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new PReluPluginDynamic(serial_data, serial_length); return new PReluPluginDynamic(serial_data, serial_length);
} }
}; };
......
paddle/fluid/inference/tensorrt/plugin/qkv_to_context_plugin.cu
浏览文件 @ b5aab4f0
...@@ -147,11 +147,11 @@ inline void TransposeQKV(const int batch, const int seq_len, ...@@ -147,11 +147,11 @@ inline void TransposeQKV(const int batch, const int seq_len,
} }
} }
int QkvToContextPluginDynamic::initialize() { return 0; } int QkvToContextPluginDynamic::initialize() TRT_NOEXCEPT { return 0; }
nvinfer1::DimsExprs QkvToContextPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs QkvToContextPluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs,
nvinfer1::IExprBuilder &expr_builder) { nvinfer1::IExprBuilder &expr_builder) TRT_NOEXCEPT {
// input[0], (B, S, 3 * N * H, 1, 1) // input[0], (B, S, 3 * N * H, 1, 1)
// input[1], (B, head_num, seq_len, seq_len) // input[1], (B, head_num, seq_len, seq_len)
// output, (B, seq_len, hidden) // output, (B, seq_len, hidden)
...@@ -177,7 +177,7 @@ nvinfer1::DimsExprs QkvToContextPluginDynamic::getOutputDimensions( ...@@ -177,7 +177,7 @@ nvinfer1::DimsExprs QkvToContextPluginDynamic::getOutputDimensions(
bool QkvToContextPluginDynamic::supportsFormatCombination( bool QkvToContextPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of swish plugin shoule not be nullptr.")); "The input of swish plugin shoule not be nullptr."));
...@@ -215,7 +215,8 @@ bool QkvToContextPluginDynamic::supportsFormatCombination( ...@@ -215,7 +215,8 @@ bool QkvToContextPluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType QkvToContextPluginDynamic::getOutputDataType( nvinfer1::DataType QkvToContextPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType *input_types, int nb_inputs) const { int index, const nvinfer1::DataType *input_types,
int nb_inputs) const TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ( PADDLE_ENFORCE_EQ(
index, 0, platform::errors::InvalidArgument( index, 0, platform::errors::InvalidArgument(
"The EmbEltwiseLayernorm Plugin only has one input, so the " "The EmbEltwiseLayernorm Plugin only has one input, so the "
...@@ -235,7 +236,7 @@ __global__ void apply_scale(T *data, T scale, int n) { ...@@ -235,7 +236,7 @@ __global__ void apply_scale(T *data, T scale, int n) {
int QkvToContextPluginDynamic::enqueue( int QkvToContextPluginDynamic::enqueue(
const nvinfer1::PluginTensorDesc *input_desc, const nvinfer1::PluginTensorDesc *input_desc,
const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs, const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs,
void *const *outputs, void *workspace, cudaStream_t stream) { void *const *outputs, void *workspace, cudaStream_t stream) TRT_NOEXCEPT {
auto input_dims = input_desc[0].dims; auto input_dims = input_desc[0].dims;
int input_num = ProductDim(input_dims); int input_num = ProductDim(input_dims);
// input[0], (B, S, 3 * N * H, 1, 1) // input[0], (B, S, 3 * N * H, 1, 1)
......
paddle/fluid/inference/tensorrt/plugin/qkv_to_context_plugin.h
浏览文件 @ b5aab4f0
...@@ -59,21 +59,23 @@ class QkvToContextPluginDynamic : public DynamicPluginTensorRT { ...@@ -59,21 +59,23 @@ class QkvToContextPluginDynamic : public DynamicPluginTensorRT {
DeserializeValue(&serial_data, &serial_length, &scale_); DeserializeValue(&serial_data, &serial_length, &scale_);
DeserializeValue(&serial_data, &serial_length, &with_fp16_); DeserializeValue(&serial_data, &serial_length, &with_fp16_);
} }
nvinfer1::IPluginV2DynamicExt* clone() const override { nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override {
return new QkvToContextPluginDynamic(hidden_, head_number_, head_size_, return new QkvToContextPluginDynamic(hidden_, head_number_, head_size_,
scale_, with_fp16_); scale_, with_fp16_);
} }
const char* getPluginType() const override { return "qkv_to_context_plugin"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "qkv_to_context_plugin";
int initialize() override; }
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
int initialize() TRT_NOEXCEPT override;
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
return SerializedSize(hidden_) + SerializedSize(head_number_) + return SerializedSize(hidden_) + SerializedSize(head_number_) +
SerializedSize(head_size_) + SerializedSize(scale_) + SerializedSize(head_size_) + SerializedSize(scale_) +
SerializedSize(with_fp16_); SerializedSize(with_fp16_);
} }
void serialize(void* buffer) const override { void serialize(void* buffer) const TRT_NOEXCEPT override {
SerializeValue(&buffer, hidden_); SerializeValue(&buffer, hidden_);
SerializeValue(&buffer, head_number_); SerializeValue(&buffer, head_number_);
SerializeValue(&buffer, head_size_); SerializeValue(&buffer, head_size_);
...@@ -83,33 +85,34 @@ class QkvToContextPluginDynamic : public DynamicPluginTensorRT { ...@@ -83,33 +85,34 @@ class QkvToContextPluginDynamic : public DynamicPluginTensorRT {
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) override; nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* in_out, const nvinfer1::PluginTensorDesc* in_out,
int nb_inputs, int nb_outputs) override; int nb_inputs,
int nb_outputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nb_inputs, int nb_inputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nb_outputs) override {} int nb_outputs) TRT_NOEXCEPT override {}
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nb_inputs, int nb_inputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nb_outputs) const override { int nb_outputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* input_types, int index, const nvinfer1::DataType* input_types,
int nb_inputs) const override; int nb_inputs) const TRT_NOEXCEPT override;
void destroy() override { delete this; } void destroy() TRT_NOEXCEPT override { delete this; }
private: private:
int hidden_; int hidden_;
...@@ -121,31 +124,34 @@ class QkvToContextPluginDynamic : public DynamicPluginTensorRT { ...@@ -121,31 +124,34 @@ class QkvToContextPluginDynamic : public DynamicPluginTensorRT {
class QkvToContextPluginDynamicCreator : public nvinfer1::IPluginCreator { class QkvToContextPluginDynamicCreator : public nvinfer1::IPluginCreator {
public: public:
QkvToContextPluginDynamicCreator() {} QkvToContextPluginDynamicCreator() {}
const char* getPluginName() const override { return "qkv_to_context_plugin"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "qkv_to_context_plugin";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
const nvinfer1::PluginFieldCollection* getFieldNames() override { const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2* createPlugin( nvinfer1::IPluginV2* createPlugin(const char* name,
const char* name, const nvinfer1::PluginFieldCollection* fc) override { const nvinfer1::PluginFieldCollection* fc)
TRT_NOEXCEPT override {
return nullptr; return nullptr;
} }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
auto plugin = new QkvToContextPluginDynamic(serial_data, serial_length); auto plugin = new QkvToContextPluginDynamic(serial_data, serial_length);
return plugin; return plugin;
} }
void setPluginNamespace(const char* lib_namespace) override { void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override {
plugin_namespace_ = lib_namespace; plugin_namespace_ = lib_namespace;
} }
const char* getPluginNamespace() const override { const char* getPluginNamespace() const TRT_NOEXCEPT override {
return plugin_namespace_.c_str(); return plugin_namespace_.c_str();
} }
......
paddle/fluid/inference/tensorrt/plugin/roi_align_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -200,7 +200,8 @@ RoiAlignPluginDynamic::RoiAlignPluginDynamic(void const* data, size_t length) { ...@@ -200,7 +200,8 @@ RoiAlignPluginDynamic::RoiAlignPluginDynamic(void const* data, size_t length) {
smem_per_block_ = smem_per_block; smem_per_block_ = smem_per_block;
} }
nvinfer1::IPluginV2DynamicExt* RoiAlignPluginDynamic::clone() const { nvinfer1::IPluginV2DynamicExt* RoiAlignPluginDynamic::clone() const
TRT_NOEXCEPT {
auto* plugin = auto* plugin =
new RoiAlignPluginDynamic(data_type_, pooled_height_, pooled_width_, new RoiAlignPluginDynamic(data_type_, pooled_height_, pooled_width_,
spatial_scale_, sampling_ratio_); spatial_scale_, sampling_ratio_);
...@@ -210,7 +211,7 @@ nvinfer1::IPluginV2DynamicExt* RoiAlignPluginDynamic::clone() const { ...@@ -210,7 +211,7 @@ nvinfer1::IPluginV2DynamicExt* RoiAlignPluginDynamic::clone() const {
nvinfer1::DimsExprs RoiAlignPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs RoiAlignPluginDynamic::getOutputDimensions(
int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs, int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs,
nvinfer1::IExprBuilder& exprBuilder) { nvinfer1::IExprBuilder& exprBuilder) TRT_NOEXCEPT {
nvinfer1::DimsExprs ret{}; nvinfer1::DimsExprs ret{};
ret.nbDims = 4; ret.nbDims = 4;
ret.d[0] = inputs[1].d[0]; // roi ret.d[0] = inputs[1].d[0]; // roi
...@@ -222,7 +223,7 @@ nvinfer1::DimsExprs RoiAlignPluginDynamic::getOutputDimensions( ...@@ -222,7 +223,7 @@ nvinfer1::DimsExprs RoiAlignPluginDynamic::getOutputDimensions(
bool RoiAlignPluginDynamic::supportsFormatCombination( bool RoiAlignPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc* inOut, int nbInputs, int pos, const nvinfer1::PluginTensorDesc* inOut, int nbInputs,
int nbOutputs) { int nbOutputs) TRT_NOEXCEPT {
if (inOut[pos].format != nvinfer1::TensorFormat::kLINEAR) { if (inOut[pos].format != nvinfer1::TensorFormat::kLINEAR) {
return false; return false;
} }
...@@ -234,11 +235,12 @@ bool RoiAlignPluginDynamic::supportsFormatCombination( ...@@ -234,11 +235,12 @@ bool RoiAlignPluginDynamic::supportsFormatCombination(
void RoiAlignPluginDynamic::configurePlugin( void RoiAlignPluginDynamic::configurePlugin(
const nvinfer1::DynamicPluginTensorDesc* in, int nbInputs, const nvinfer1::DynamicPluginTensorDesc* in, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, int nbOutputs) {} const nvinfer1::DynamicPluginTensorDesc* out, int nbOutputs) TRT_NOEXCEPT {}
size_t RoiAlignPluginDynamic::getWorkspaceSize( size_t RoiAlignPluginDynamic::getWorkspaceSize(
const nvinfer1::PluginTensorDesc* inputs, int nbInputs, const nvinfer1::PluginTensorDesc* inputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, int nbOutputs) const { const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const TRT_NOEXCEPT {
return 0; return 0;
} }
...@@ -287,7 +289,7 @@ int RoiAlignPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc* inputDesc, ...@@ -287,7 +289,7 @@ int RoiAlignPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, const void* const* inputs,
void* const* outputs, void* workspace, void* const* outputs, void* workspace,
cudaStream_t stream) { cudaStream_t stream) TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(outputDesc[0].type, data_type_, PADDLE_ENFORCE_EQ(outputDesc[0].type, data_type_,
platform::errors::InvalidArgument( platform::errors::InvalidArgument(
"TRT RoiAlignPluginDynamic expects outputDesc[0].type " "TRT RoiAlignPluginDynamic expects outputDesc[0].type "
...@@ -302,21 +304,22 @@ int RoiAlignPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc* inputDesc, ...@@ -302,21 +304,22 @@ int RoiAlignPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
} }
nvinfer1::DataType RoiAlignPluginDynamic::getOutputDataType( nvinfer1::DataType RoiAlignPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType* inputTypes, int nbInputs) const { int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const TRT_NOEXCEPT {
return inputTypes[0]; return inputTypes[0];
} }
const char* RoiAlignPluginDynamic::getPluginType() const { const char* RoiAlignPluginDynamic::getPluginType() const TRT_NOEXCEPT {
return "roi_align_plugin_dynamic"; return "roi_align_plugin_dynamic";
} }
int RoiAlignPluginDynamic::getNbOutputs() const { return 1; } int RoiAlignPluginDynamic::getNbOutputs() const TRT_NOEXCEPT { return 1; }
int RoiAlignPluginDynamic::initialize() { return 0; } int RoiAlignPluginDynamic::initialize() TRT_NOEXCEPT { return 0; }
void RoiAlignPluginDynamic::terminate() {} void RoiAlignPluginDynamic::terminate() TRT_NOEXCEPT {}
size_t RoiAlignPluginDynamic::getSerializationSize() const { size_t RoiAlignPluginDynamic::getSerializationSize() const TRT_NOEXCEPT {
size_t serialize_size = 0; size_t serialize_size = 0;
serialize_size += SerializedSize(data_type_); serialize_size += SerializedSize(data_type_);
serialize_size += SerializedSize(pooled_height_); serialize_size += SerializedSize(pooled_height_);
...@@ -326,7 +329,7 @@ size_t RoiAlignPluginDynamic::getSerializationSize() const { ...@@ -326,7 +329,7 @@ size_t RoiAlignPluginDynamic::getSerializationSize() const {
return serialize_size; return serialize_size;
} }
void RoiAlignPluginDynamic::serialize(void* buffer) const { void RoiAlignPluginDynamic::serialize(void* buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, data_type_); SerializeValue(&buffer, data_type_);
SerializeValue(&buffer, pooled_height_); SerializeValue(&buffer, pooled_height_);
SerializeValue(&buffer, pooled_width_); SerializeValue(&buffer, pooled_width_);
...@@ -334,40 +337,43 @@ void RoiAlignPluginDynamic::serialize(void* buffer) const { ...@@ -334,40 +337,43 @@ void RoiAlignPluginDynamic::serialize(void* buffer) const {
SerializeValue(&buffer, sampling_ratio_); SerializeValue(&buffer, sampling_ratio_);
} }
void RoiAlignPluginDynamic::destroy() {} void RoiAlignPluginDynamic::destroy() TRT_NOEXCEPT {}
RoiAlignPluginDynamicCreator::RoiAlignPluginDynamicCreator() {} RoiAlignPluginDynamicCreator::RoiAlignPluginDynamicCreator() {}
void RoiAlignPluginDynamicCreator::setPluginNamespace( void RoiAlignPluginDynamicCreator::setPluginNamespace(const char* lib_namespace)
const char* lib_namespace) { TRT_NOEXCEPT {
namespace_ = std::string(lib_namespace); namespace_ = std::string(lib_namespace);
} }
const char* RoiAlignPluginDynamicCreator::getPluginNamespace() const { const char* RoiAlignPluginDynamicCreator::getPluginNamespace() const
TRT_NOEXCEPT {
return namespace_.c_str(); return namespace_.c_str();
} }
const char* RoiAlignPluginDynamicCreator::getPluginName() const { const char* RoiAlignPluginDynamicCreator::getPluginName() const TRT_NOEXCEPT {
return "roi_align_plugin_dynamic"; return "roi_align_plugin_dynamic";
} }
const char* RoiAlignPluginDynamicCreator::getPluginVersion() const { const char* RoiAlignPluginDynamicCreator::getPluginVersion() const
TRT_NOEXCEPT {
return "1"; return "1";
} }
const nvinfer1::PluginFieldCollection* const nvinfer1::PluginFieldCollection*
RoiAlignPluginDynamicCreator::getFieldNames() { RoiAlignPluginDynamicCreator::getFieldNames() TRT_NOEXCEPT {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2Ext* RoiAlignPluginDynamicCreator::createPlugin( nvinfer1::IPluginV2Ext* RoiAlignPluginDynamicCreator::createPlugin(
const char* name, const nvinfer1::PluginFieldCollection* fc) { const char* name, const nvinfer1::PluginFieldCollection* fc) TRT_NOEXCEPT {
const nvinfer1::PluginField* fields = fc->fields; const nvinfer1::PluginField* fields = fc->fields;
return nullptr; return nullptr;
} }
nvinfer1::IPluginV2Ext* RoiAlignPluginDynamicCreator::deserializePlugin( nvinfer1::IPluginV2Ext* RoiAlignPluginDynamicCreator::deserializePlugin(
const char* name, const void* serial_data, size_t serial_length) { const char* name, const void* serial_data,
size_t serial_length) TRT_NOEXCEPT {
auto plugin = new RoiAlignPluginDynamic(serial_data, serial_length); auto plugin = new RoiAlignPluginDynamic(serial_data, serial_length);
plugin->setPluginNamespace(namespace_.c_str()); plugin->setPluginNamespace(namespace_.c_str());
return plugin; return plugin;
......
paddle/fluid/inference/tensorrt/plugin/roi_align_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -34,37 +34,38 @@ class RoiAlignPluginDynamic : public DynamicPluginTensorRT { ...@@ -34,37 +34,38 @@ class RoiAlignPluginDynamic : public DynamicPluginTensorRT {
int sampling_ratio); int sampling_ratio);
RoiAlignPluginDynamic(void const* data, size_t length); RoiAlignPluginDynamic(void const* data, size_t length);
~RoiAlignPluginDynamic() = default; ~RoiAlignPluginDynamic() = default;
nvinfer1::IPluginV2DynamicExt* clone() const override; nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override;
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs, int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs,
nvinfer1::IExprBuilder& exprBuilder) override; nvinfer1::IExprBuilder& exprBuilder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* inOut, const nvinfer1::PluginTensorDesc* inOut,
int nbInputs, int nbOutputs) override; int nbInputs,
int nbOutputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nbInputs, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nbOutputs) override; int nbOutputs) TRT_NOEXCEPT override;
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nbInputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const override; int nbOutputs) const TRT_NOEXCEPT override;
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* inputTypes, int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const override; int nbInputs) const TRT_NOEXCEPT override;
const char* getPluginType() const override; const char* getPluginType() const TRT_NOEXCEPT override;
int getNbOutputs() const override; int getNbOutputs() const TRT_NOEXCEPT override;
int initialize() override; int initialize() TRT_NOEXCEPT override;
void terminate() override; void terminate() TRT_NOEXCEPT override;
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
void destroy() override; void destroy() TRT_NOEXCEPT override;
private: private:
template <typename T, typename OutT> template <typename T, typename OutT>
...@@ -87,17 +88,18 @@ class RoiAlignPluginDynamicCreator : public nvinfer1::IPluginCreator { ...@@ -87,17 +88,18 @@ class RoiAlignPluginDynamicCreator : public nvinfer1::IPluginCreator {
RoiAlignPluginDynamicCreator(); RoiAlignPluginDynamicCreator();
~RoiAlignPluginDynamicCreator() override = default; ~RoiAlignPluginDynamicCreator() override = default;
void setPluginNamespace(const char* lib_namespace) override; void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override;
const char* getPluginNamespace() const override; const char* getPluginNamespace() const TRT_NOEXCEPT override;
const char* getPluginName() const override; const char* getPluginName() const TRT_NOEXCEPT override;
const char* getPluginVersion() const override; const char* getPluginVersion() const TRT_NOEXCEPT override;
const nvinfer1::PluginFieldCollection* getFieldNames() override; const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override;
nvinfer1::IPluginV2Ext* createPlugin( nvinfer1::IPluginV2Ext* createPlugin(
const char* name, const nvinfer1::PluginFieldCollection* fc) override; const char* name,
nvinfer1::IPluginV2Ext* deserializePlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) TRT_NOEXCEPT override;
const void* serial_data, nvinfer1::IPluginV2Ext* deserializePlugin(
size_t serial_length) override; const char* name, const void* serial_data,
size_t serial_length) TRT_NOEXCEPT override;
private: private:
std::string namespace_; std::string namespace_;
......
paddle/fluid/inference/tensorrt/plugin/skip_layernorm_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -29,7 +29,7 @@ namespace plugin { ...@@ -29,7 +29,7 @@ namespace plugin {
// Dynamic Plugin below. // Dynamic Plugin below.
#if IS_TRT_VERSION_GE(6000) #if IS_TRT_VERSION_GE(6000)
int SkipLayerNormPluginDynamic::initialize() { int SkipLayerNormPluginDynamic::initialize() TRT_NOEXCEPT {
cudaMalloc(&bias_gpu_, sizeof(float) * bias_size_); cudaMalloc(&bias_gpu_, sizeof(float) * bias_size_);
cudaMemcpy(bias_gpu_, bias_.data(), bias_size_ * sizeof(float), cudaMemcpy(bias_gpu_, bias_.data(), bias_size_ * sizeof(float),
cudaMemcpyHostToDevice); cudaMemcpyHostToDevice);
...@@ -39,7 +39,7 @@ int SkipLayerNormPluginDynamic::initialize() { ...@@ -39,7 +39,7 @@ int SkipLayerNormPluginDynamic::initialize() {
return 0; return 0;
} }
void SkipLayerNormPluginDynamic::terminate() { void SkipLayerNormPluginDynamic::terminate() TRT_NOEXCEPT {
if (bias_gpu_) { if (bias_gpu_) {
cudaFree(bias_gpu_); cudaFree(bias_gpu_);
bias_gpu_ = nullptr; bias_gpu_ = nullptr;
...@@ -52,13 +52,13 @@ void SkipLayerNormPluginDynamic::terminate() { ...@@ -52,13 +52,13 @@ void SkipLayerNormPluginDynamic::terminate() {
nvinfer1::DimsExprs SkipLayerNormPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs SkipLayerNormPluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs,
nvinfer1::IExprBuilder &expr_builder) { nvinfer1::IExprBuilder &expr_builder) TRT_NOEXCEPT {
return inputs[0]; return inputs[0];
} }
bool SkipLayerNormPluginDynamic::supportsFormatCombination( bool SkipLayerNormPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of swish plugin shoule not be nullptr.")); "The input of swish plugin shoule not be nullptr."));
...@@ -96,7 +96,8 @@ bool SkipLayerNormPluginDynamic::supportsFormatCombination( ...@@ -96,7 +96,8 @@ bool SkipLayerNormPluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType SkipLayerNormPluginDynamic::getOutputDataType( nvinfer1::DataType SkipLayerNormPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType *input_types, int nb_inputs) const { int index, const nvinfer1::DataType *input_types,
int nb_inputs) const TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(index, 0, PADDLE_ENFORCE_EQ(index, 0,
platform::errors::InvalidArgument( platform::errors::InvalidArgument(
"The SkipLayerNorm Plugin only has one input, so the " "The SkipLayerNorm Plugin only has one input, so the "
...@@ -112,7 +113,7 @@ nvinfer1::DataType SkipLayerNormPluginDynamic::getOutputDataType( ...@@ -112,7 +113,7 @@ nvinfer1::DataType SkipLayerNormPluginDynamic::getOutputDataType(
int SkipLayerNormPluginDynamic::enqueue( int SkipLayerNormPluginDynamic::enqueue(
const nvinfer1::PluginTensorDesc *input_desc, const nvinfer1::PluginTensorDesc *input_desc,
const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs, const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs,
void *const *outputs, void *workspace, cudaStream_t stream) { void *const *outputs, void *workspace, cudaStream_t stream) TRT_NOEXCEPT {
auto input_dims = input_desc[0].dims; auto input_dims = input_desc[0].dims;
size_t num = ProductDim(input_dims); size_t num = ProductDim(input_dims);
int hidden = input_dims.d[2]; int hidden = input_dims.d[2];
......
paddle/fluid/inference/tensorrt/plugin/skip_layernorm_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -39,6 +39,7 @@ class SkipLayerNormPluginDynamic : public DynamicPluginTensorRT { ...@@ -39,6 +39,7 @@ class SkipLayerNormPluginDynamic : public DynamicPluginTensorRT {
std::copy(bias, bias + bias_size, bias_.data()); std::copy(bias, bias + bias_size, bias_.data());
std::copy(scale, scale + scale_size, scale_.data()); std::copy(scale, scale + scale_size, scale_.data());
} }
SkipLayerNormPluginDynamic(void const* serial_data, size_t serial_length) { SkipLayerNormPluginDynamic(void const* serial_data, size_t serial_length) {
DeserializeValue(&serial_data, &serial_length, &bias_); DeserializeValue(&serial_data, &serial_length, &bias_);
DeserializeValue(&serial_data, &serial_length, &scale_); DeserializeValue(&serial_data, &serial_length, &scale_);
...@@ -48,7 +49,7 @@ class SkipLayerNormPluginDynamic : public DynamicPluginTensorRT { ...@@ -48,7 +49,7 @@ class SkipLayerNormPluginDynamic : public DynamicPluginTensorRT {
DeserializeValue(&serial_data, &serial_length, &with_fp16_); DeserializeValue(&serial_data, &serial_length, &with_fp16_);
} }
nvinfer1::IPluginV2DynamicExt* clone() const override { nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override {
auto ptr = new SkipLayerNormPluginDynamic( auto ptr = new SkipLayerNormPluginDynamic(
bias_.data(), scale_.data(), bias_size_, scale_size_, eps_, with_fp16_); bias_.data(), scale_.data(), bias_size_, scale_size_, eps_, with_fp16_);
ptr->bias_gpu_ = bias_gpu_; ptr->bias_gpu_ = bias_gpu_;
...@@ -56,17 +57,19 @@ class SkipLayerNormPluginDynamic : public DynamicPluginTensorRT { ...@@ -56,17 +57,19 @@ class SkipLayerNormPluginDynamic : public DynamicPluginTensorRT {
return ptr; return ptr;
} }
const char* getPluginType() const override { return "skip_layernorm_plugin"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "skip_layernorm_plugin";
int initialize() override; }
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
int initialize() TRT_NOEXCEPT override;
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
size_t ser_size = SerializedSize(bias_) + SerializedSize(scale_) + size_t ser_size = SerializedSize(bias_) + SerializedSize(scale_) +
SerializedSize(bias_size_) + SerializedSize(scale_size_) + SerializedSize(bias_size_) + SerializedSize(scale_size_) +
SerializedSize(eps_) + SerializedSize(with_fp16_); SerializedSize(eps_) + SerializedSize(with_fp16_);
return ser_size; return ser_size;
} }
void serialize(void* buffer) const override { void serialize(void* buffer) const TRT_NOEXCEPT override {
SerializeValue(&buffer, bias_); SerializeValue(&buffer, bias_);
SerializeValue(&buffer, scale_); SerializeValue(&buffer, scale_);
SerializeValue(&buffer, bias_size_); SerializeValue(&buffer, bias_size_);
...@@ -77,34 +80,35 @@ class SkipLayerNormPluginDynamic : public DynamicPluginTensorRT { ...@@ -77,34 +80,35 @@ class SkipLayerNormPluginDynamic : public DynamicPluginTensorRT {
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) override; nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* in_out, const nvinfer1::PluginTensorDesc* in_out,
int nb_inputs, int nb_outputs) override; int nb_inputs,
int nb_outputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nb_inputs, int nb_inputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nb_outputs) override {} int nb_outputs) TRT_NOEXCEPT override {}
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nb_inputs, int nb_inputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nb_outputs) const override { int nb_outputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
int enqueue(const nvinfer1::PluginTensorDesc* input_desc, int enqueue(const nvinfer1::PluginTensorDesc* input_desc,
const nvinfer1::PluginTensorDesc* output_desc, const nvinfer1::PluginTensorDesc* output_desc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* input_types, int index, const nvinfer1::DataType* input_types,
int nb_inputs) const override; int nb_inputs) const TRT_NOEXCEPT override;
void destroy() override { delete this; } void destroy() TRT_NOEXCEPT override { delete this; }
void terminate() override; void terminate() TRT_NOEXCEPT override;
private: private:
std::vector<float> bias_; std::vector<float> bias_;
...@@ -122,31 +126,34 @@ class SkipLayerNormPluginDynamic : public DynamicPluginTensorRT { ...@@ -122,31 +126,34 @@ class SkipLayerNormPluginDynamic : public DynamicPluginTensorRT {
class SkipLayerNormPluginDynamicCreator : public nvinfer1::IPluginCreator { class SkipLayerNormPluginDynamicCreator : public nvinfer1::IPluginCreator {
public: public:
SkipLayerNormPluginDynamicCreator() {} SkipLayerNormPluginDynamicCreator() {}
const char* getPluginName() const override { return "skip_layernorm_plugin"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "skip_layernorm_plugin";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
const nvinfer1::PluginFieldCollection* getFieldNames() override { const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2* createPlugin( nvinfer1::IPluginV2* createPlugin(const char* name,
const char* name, const nvinfer1::PluginFieldCollection* fc) override { const nvinfer1::PluginFieldCollection* fc)
TRT_NOEXCEPT override {
return nullptr; return nullptr;
} }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
auto plugin = new SkipLayerNormPluginDynamic(serial_data, serial_length); auto plugin = new SkipLayerNormPluginDynamic(serial_data, serial_length);
return plugin; return plugin;
} }
void setPluginNamespace(const char* lib_namespace) override { void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override {
plugin_namespace_ = lib_namespace; plugin_namespace_ = lib_namespace;
} }
const char* getPluginNamespace() const override { const char* getPluginNamespace() const TRT_NOEXCEPT override {
return plugin_namespace_.c_str(); return plugin_namespace_.c_str();
} }
......
paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -75,12 +75,12 @@ SlicePlugin::~SlicePlugin() { ...@@ -75,12 +75,12 @@ SlicePlugin::~SlicePlugin() {
cudaFree(offset_temp_data_); cudaFree(offset_temp_data_);
} }
SlicePlugin *SlicePlugin::clone() const { SlicePlugin *SlicePlugin::clone() const TRT_NOEXCEPT {
return new SlicePlugin(starts_, ends_, axes_, with_fp16_); return new SlicePlugin(starts_, ends_, axes_, with_fp16_);
} }
bool SlicePlugin::supportsFormat(nvinfer1::DataType type, bool SlicePlugin::supportsFormat(
nvinfer1::PluginFormat format) const { nvinfer1::DataType type, nvinfer1::PluginFormat format) const TRT_NOEXCEPT {
if (with_fp16_) { if (with_fp16_) {
return ((type == nvinfer1::DataType::kFLOAT || return ((type == nvinfer1::DataType::kFLOAT ||
type == nvinfer1::DataType::kHALF) && type == nvinfer1::DataType::kHALF) &&
...@@ -91,9 +91,8 @@ bool SlicePlugin::supportsFormat(nvinfer1::DataType type, ...@@ -91,9 +91,8 @@ bool SlicePlugin::supportsFormat(nvinfer1::DataType type,
} }
} }
nvinfer1::Dims SlicePlugin::getOutputDimensions(int index, nvinfer1::Dims SlicePlugin::getOutputDimensions(
const nvinfer1::Dims *inputs, int index, const nvinfer1::Dims *inputs, int nb_input_dims) TRT_NOEXCEPT {
int nb_input_dims) {
auto in_dims = inputs[0]; auto in_dims = inputs[0];
nvinfer1::Dims out_dims = in_dims; nvinfer1::Dims out_dims = in_dims;
for (size_t i = 0; i < axes_.size(); i++) { for (size_t i = 0; i < axes_.size(); i++) {
...@@ -109,7 +108,7 @@ int SlicePlugin::enqueue(int batch_size, const void *const *inputs, ...@@ -109,7 +108,7 @@ int SlicePlugin::enqueue(int batch_size, const void *const *inputs,
void **outputs, void *workspace, cudaStream_t stream) { void **outputs, void *workspace, cudaStream_t stream) {
#else #else
void *const *outputs, void *workspace, void *const *outputs, void *workspace,
cudaStream_t stream) { cudaStream_t stream) TRT_NOEXCEPT {
#endif #endif
auto input_dims = getInputDims(0); auto input_dims = getInputDims(0);
...@@ -187,13 +186,13 @@ int SlicePlugin::enqueue(int batch_size, const void *const *inputs, ...@@ -187,13 +186,13 @@ int SlicePlugin::enqueue(int batch_size, const void *const *inputs,
return cudaGetLastError() != cudaSuccess; return cudaGetLastError() != cudaSuccess;
} }
size_t SlicePlugin::getSerializationSize() const { size_t SlicePlugin::getSerializationSize() const TRT_NOEXCEPT {
return getBaseSerializationSize() + SerializedSize(getPluginType()) + return getBaseSerializationSize() + SerializedSize(getPluginType()) +
SerializedSize(starts_) + SerializedSize(ends_) + SerializedSize(starts_) + SerializedSize(ends_) +
SerializedSize(axes_); SerializedSize(axes_);
} }
void SlicePlugin::serialize(void *buffer) const { void SlicePlugin::serialize(void *buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, getPluginType()); SerializeValue(&buffer, getPluginType());
serializeBase(buffer); serializeBase(buffer);
SerializeValue(&buffer, starts_); SerializeValue(&buffer, starts_);
...@@ -222,23 +221,23 @@ SlicePluginDynamic::SlicePluginDynamic(void const *serialData, ...@@ -222,23 +221,23 @@ SlicePluginDynamic::SlicePluginDynamic(void const *serialData,
cudaStreamCreate(&copy_stream_); cudaStreamCreate(&copy_stream_);
} }
void SlicePluginDynamic::destroy() { void SlicePluginDynamic::destroy() TRT_NOEXCEPT {
cudaStreamDestroy(copy_stream_); cudaStreamDestroy(copy_stream_);
cudaEventDestroy(copy_event_); cudaEventDestroy(copy_event_);
cudaFree(offset_temp_data_); cudaFree(offset_temp_data_);
delete this; delete this;
} }
int SlicePluginDynamic::initialize() { return 0; } int SlicePluginDynamic::initialize() TRT_NOEXCEPT { return 0; }
size_t SlicePluginDynamic::getSerializationSize() const { size_t SlicePluginDynamic::getSerializationSize() const TRT_NOEXCEPT {
size_t size = SerializedSize(starts_) + SerializedSize(ends_) + size_t size = SerializedSize(starts_) + SerializedSize(ends_) +
SerializedSize(axes_) + SerializedSize(with_fp16_); SerializedSize(axes_) + SerializedSize(with_fp16_);
return size; return size;
} }
void SlicePluginDynamic::serialize(void *buffer) const { void SlicePluginDynamic::serialize(void *buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, starts_); SerializeValue(&buffer, starts_);
SerializeValue(&buffer, ends_); SerializeValue(&buffer, ends_);
SerializeValue(&buffer, axes_); SerializeValue(&buffer, axes_);
...@@ -247,7 +246,7 @@ void SlicePluginDynamic::serialize(void *buffer) const { ...@@ -247,7 +246,7 @@ void SlicePluginDynamic::serialize(void *buffer) const {
nvinfer1::DimsExprs SlicePluginDynamic::getOutputDimensions( nvinfer1::DimsExprs SlicePluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs,
nvinfer1::IExprBuilder &expr_builder) { nvinfer1::IExprBuilder &expr_builder) TRT_NOEXCEPT {
auto in_dims = inputs[0]; auto in_dims = inputs[0];
nvinfer1::DimsExprs ret = in_dims; nvinfer1::DimsExprs ret = in_dims;
// start, ends should greater 0 // start, ends should greater 0
...@@ -261,7 +260,7 @@ nvinfer1::DimsExprs SlicePluginDynamic::getOutputDimensions( ...@@ -261,7 +260,7 @@ nvinfer1::DimsExprs SlicePluginDynamic::getOutputDimensions(
bool SlicePluginDynamic::supportsFormatCombination( bool SlicePluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of swish plugin shoule not be nullptr.")); "The input of swish plugin shoule not be nullptr."));
...@@ -289,7 +288,8 @@ bool SlicePluginDynamic::supportsFormatCombination( ...@@ -289,7 +288,8 @@ bool SlicePluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType SlicePluginDynamic::getOutputDataType( nvinfer1::DataType SlicePluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType *input_types, int nb_inputs) const { int index, const nvinfer1::DataType *input_types,
int nb_inputs) const TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument( PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument(
"The Slice Plugin only has one input, so the " "The Slice Plugin only has one input, so the "
"index value should be 0, but get %d.", "index value should be 0, but get %d.",
...@@ -304,7 +304,8 @@ nvinfer1::DataType SlicePluginDynamic::getOutputDataType( ...@@ -304,7 +304,8 @@ nvinfer1::DataType SlicePluginDynamic::getOutputDataType(
int SlicePluginDynamic::enqueue(const nvinfer1::PluginTensorDesc *input_desc, int SlicePluginDynamic::enqueue(const nvinfer1::PluginTensorDesc *input_desc,
const nvinfer1::PluginTensorDesc *output_desc, const nvinfer1::PluginTensorDesc *output_desc,
const void *const *inputs, void *const *outputs, const void *const *inputs, void *const *outputs,
void *workspace, cudaStream_t stream) { void *workspace,
cudaStream_t stream) TRT_NOEXCEPT {
auto input_dims = input_desc[0].dims; auto input_dims = input_desc[0].dims;
auto out_dims = output_desc[0].dims; auto out_dims = output_desc[0].dims;
auto num_dims = input_dims.nbDims; auto num_dims = input_dims.nbDims;
......
paddle/fluid/inference/tensorrt/plugin/slice_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -35,27 +35,29 @@ class SlicePlugin : public PluginTensorRT { ...@@ -35,27 +35,29 @@ class SlicePlugin : public PluginTensorRT {
// It should not be called by users. // It should not be called by users.
SlicePlugin(void const* serial_data, size_t serial_length); SlicePlugin(void const* serial_data, size_t serial_length);
~SlicePlugin(); ~SlicePlugin();
SlicePlugin* clone() const override; SlicePlugin* clone() const TRT_NOEXCEPT override;
const char* getPluginType() const override { return "slice_plugin"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "slice_plugin";
int initialize() override { return 0; } }
bool supportsFormat(nvinfer1::DataType type, int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
nvinfer1::PluginFormat format) const override; int initialize() TRT_NOEXCEPT override { return 0; }
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format)
const TRT_NOEXCEPT override;
nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs,
int nb_input_dims) override; int nb_input_dims) TRT_NOEXCEPT override;
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
int enqueue(int batch_size, const void* const* inputs, void** outputs, int enqueue(int batch_size, const void* const* inputs, void** outputs,
#else #else
int enqueue(int batch_size, const void* const* inputs, void* const* outputs, int enqueue(int batch_size, const void* const* inputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream) override; void* workspace, cudaStream_t stream) TRT_NOEXCEPT override;
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
// TRT will call this func to serialize the configuration of TRT // TRT will call this func to serialize the configuration of TRT
// It should not be called by users. // It should not be called by users.
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
private: private:
std::vector<int> starts_; std::vector<int> starts_;
...@@ -68,13 +70,15 @@ class SlicePlugin : public PluginTensorRT { ...@@ -68,13 +70,15 @@ class SlicePlugin : public PluginTensorRT {
class SlicePluginCreator : public TensorRTPluginCreator { class SlicePluginCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "slice_plugin"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "slice_plugin";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new SlicePlugin(serial_data, serial_length); return new SlicePlugin(serial_data, serial_length);
} }
}; };
...@@ -86,48 +90,51 @@ class SlicePluginDynamic : public DynamicPluginTensorRT { ...@@ -86,48 +90,51 @@ class SlicePluginDynamic : public DynamicPluginTensorRT {
explicit SlicePluginDynamic(std::vector<int> starts, std::vector<int> ends, explicit SlicePluginDynamic(std::vector<int> starts, std::vector<int> ends,
std::vector<int> axes, bool with_fp16); std::vector<int> axes, bool with_fp16);
nvinfer1::IPluginV2DynamicExt* clone() const override { nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override {
return new SlicePluginDynamic(starts_, ends_, axes_, with_fp16_); return new SlicePluginDynamic(starts_, ends_, axes_, with_fp16_);
} }
SlicePluginDynamic(void const* serialData, size_t serialLength); SlicePluginDynamic(void const* serialData, size_t serialLength);
const char* getPluginType() const override { return "slice_plugin_dynamic"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "slice_plugin_dynamic";
int initialize() override; }
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
int initialize() TRT_NOEXCEPT override;
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) override; nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* inOut, const nvinfer1::PluginTensorDesc* inOut,
int nbInputs, int nbOutputs) override; int nbInputs,
int nbOutputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nbInputs, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nbOutputs) override {} int nbOutputs) TRT_NOEXCEPT override {}
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nbInputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const override { int nbOutputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* inputTypes, int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const override; int nbInputs) const TRT_NOEXCEPT override;
void destroy() override; void destroy() TRT_NOEXCEPT override;
private: private:
std::vector<int> starts_; std::vector<int> starts_;
...@@ -140,13 +147,15 @@ class SlicePluginDynamic : public DynamicPluginTensorRT { ...@@ -140,13 +147,15 @@ class SlicePluginDynamic : public DynamicPluginTensorRT {
class SlicePluginDynamicCreator : public TensorRTPluginCreator { class SlicePluginDynamicCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "slice_plugin_dynamic"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "slice_plugin_dynamic";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serialData, const char* name, const void* serialData,
size_t serialLength) override { size_t serialLength) TRT_NOEXCEPT override {
return new SlicePluginDynamic(serialData, serialLength); return new SlicePluginDynamic(serialData, serialLength);
} }
}; };
......
paddle/fluid/inference/tensorrt/plugin/special_slice_plugin.cu
浏览文件 @ b5aab4f0
...@@ -30,28 +30,29 @@ SpecialSlicePluginDynamic::SpecialSlicePluginDynamic(void const* serial_data, ...@@ -30,28 +30,29 @@ SpecialSlicePluginDynamic::SpecialSlicePluginDynamic(void const* serial_data,
SpecialSlicePluginDynamic::~SpecialSlicePluginDynamic() {} SpecialSlicePluginDynamic::~SpecialSlicePluginDynamic() {}
nvinfer1::IPluginV2DynamicExt* SpecialSlicePluginDynamic::clone() const { nvinfer1::IPluginV2DynamicExt* SpecialSlicePluginDynamic::clone() const
TRT_NOEXCEPT {
return new SpecialSlicePluginDynamic(); return new SpecialSlicePluginDynamic();
} }
const char* SpecialSlicePluginDynamic::getPluginType() const { const char* SpecialSlicePluginDynamic::getPluginType() const TRT_NOEXCEPT {
return "special_slice_plugin"; return "special_slice_plugin";
} }
int SpecialSlicePluginDynamic::getNbOutputs() const { return 1; } int SpecialSlicePluginDynamic::getNbOutputs() const TRT_NOEXCEPT { return 1; }
int SpecialSlicePluginDynamic::initialize() { return 0; } int SpecialSlicePluginDynamic::initialize() TRT_NOEXCEPT { return 0; }
size_t SpecialSlicePluginDynamic::getSerializationSize() const { size_t SpecialSlicePluginDynamic::getSerializationSize() const TRT_NOEXCEPT {
size_t serialize_size = 0; size_t serialize_size = 0;
return serialize_size; return serialize_size;
} }
void SpecialSlicePluginDynamic::serialize(void* buffer) const {} void SpecialSlicePluginDynamic::serialize(void* buffer) const TRT_NOEXCEPT {}
nvinfer1::DimsExprs SpecialSlicePluginDynamic::getOutputDimensions( nvinfer1::DimsExprs SpecialSlicePluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) { nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT {
nvinfer1::DimsExprs output(inputs[0]); nvinfer1::DimsExprs output(inputs[0]);
output.nbDims++; output.nbDims++;
for (int i = output.nbDims - 1; i > 1; i--) { for (int i = output.nbDims - 1; i > 1; i--) {
...@@ -69,21 +70,22 @@ nvinfer1::DimsExprs SpecialSlicePluginDynamic::getOutputDimensions( ...@@ -69,21 +70,22 @@ nvinfer1::DimsExprs SpecialSlicePluginDynamic::getOutputDimensions(
void SpecialSlicePluginDynamic::configurePlugin( void SpecialSlicePluginDynamic::configurePlugin(
const nvinfer1::DynamicPluginTensorDesc* in, int nbInputs, const nvinfer1::DynamicPluginTensorDesc* in, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, int nbOutputs) {} const nvinfer1::DynamicPluginTensorDesc* out, int nbOutputs) TRT_NOEXCEPT {}
size_t SpecialSlicePluginDynamic::getWorkspaceSize( size_t SpecialSlicePluginDynamic::getWorkspaceSize(
const nvinfer1::PluginTensorDesc* inputs, int nbInputs, const nvinfer1::PluginTensorDesc* inputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, int nbOutputs) const { const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const TRT_NOEXCEPT {
return 0; return 0;
} }
void SpecialSlicePluginDynamic::destroy() { delete this; } void SpecialSlicePluginDynamic::destroy() TRT_NOEXCEPT { delete this; }
void SpecialSlicePluginDynamic::terminate() {} void SpecialSlicePluginDynamic::terminate() TRT_NOEXCEPT {}
bool SpecialSlicePluginDynamic::supportsFormatCombination( bool SpecialSlicePluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc* desc, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc* desc, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
if (pos == 0) // slice tensor if (pos == 0) // slice tensor
return (desc[pos].type == nvinfer1::DataType::kHALF && return (desc[pos].type == nvinfer1::DataType::kHALF &&
desc[pos].format == desc[pos].format ==
...@@ -101,7 +103,8 @@ bool SpecialSlicePluginDynamic::supportsFormatCombination( ...@@ -101,7 +103,8 @@ bool SpecialSlicePluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType SpecialSlicePluginDynamic::getOutputDataType( nvinfer1::DataType SpecialSlicePluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType* input_types, int nb_inputs) const { int index, const nvinfer1::DataType* input_types,
int nb_inputs) const TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument( PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument(
"The index should be equal to 0")); "The index should be equal to 0"));
return input_types[0]; return input_types[0];
...@@ -120,7 +123,7 @@ __global__ void SpecialSliceKernel(const T* slice_input, ...@@ -120,7 +123,7 @@ __global__ void SpecialSliceKernel(const T* slice_input,
int SpecialSlicePluginDynamic::enqueue( int SpecialSlicePluginDynamic::enqueue(
const nvinfer1::PluginTensorDesc* input_desc, const nvinfer1::PluginTensorDesc* input_desc,
const nvinfer1::PluginTensorDesc* output_desc, const void* const* inputs, const nvinfer1::PluginTensorDesc* output_desc, const void* const* inputs,
void* const* outputs, void* workspace, cudaStream_t stream) { void* const* outputs, void* workspace, cudaStream_t stream) TRT_NOEXCEPT {
auto input_dims = input_desc[0].dims; // (sum(S), 768, 1, 1) auto input_dims = input_desc[0].dims; // (sum(S), 768, 1, 1)
auto out_dims = output_desc[0].dims; // (batch, 768, 1, 1) auto out_dims = output_desc[0].dims; // (batch, 768, 1, 1)
...@@ -142,36 +145,40 @@ int SpecialSlicePluginDynamic::enqueue( ...@@ -142,36 +145,40 @@ int SpecialSlicePluginDynamic::enqueue(
SpecialSlicePluginDynamicCreator::SpecialSlicePluginDynamicCreator() {} SpecialSlicePluginDynamicCreator::SpecialSlicePluginDynamicCreator() {}
const char* SpecialSlicePluginDynamicCreator::getPluginName() const { const char* SpecialSlicePluginDynamicCreator::getPluginName() const
TRT_NOEXCEPT {
return "special_slice_plugin"; return "special_slice_plugin";
} }
const char* SpecialSlicePluginDynamicCreator::getPluginVersion() const { const char* SpecialSlicePluginDynamicCreator::getPluginVersion() const
TRT_NOEXCEPT {
return "1"; return "1";
} }
const nvinfer1::PluginFieldCollection* const nvinfer1::PluginFieldCollection*
SpecialSlicePluginDynamicCreator::getFieldNames() { SpecialSlicePluginDynamicCreator::getFieldNames() TRT_NOEXCEPT {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2* SpecialSlicePluginDynamicCreator::createPlugin( nvinfer1::IPluginV2* SpecialSlicePluginDynamicCreator::createPlugin(
const char* name, const nvinfer1::PluginFieldCollection* fc) { const char* name, const nvinfer1::PluginFieldCollection* fc) TRT_NOEXCEPT {
return new SpecialSlicePluginDynamic(); return new SpecialSlicePluginDynamic();
} }
nvinfer1::IPluginV2* SpecialSlicePluginDynamicCreator::deserializePlugin( nvinfer1::IPluginV2* SpecialSlicePluginDynamicCreator::deserializePlugin(
const char* name, const void* serial_data, size_t serial_length) { const char* name, const void* serial_data,
size_t serial_length) TRT_NOEXCEPT {
auto plugin = new SpecialSlicePluginDynamic(serial_data, serial_length); auto plugin = new SpecialSlicePluginDynamic(serial_data, serial_length);
return plugin; return plugin;
} }
void SpecialSlicePluginDynamicCreator::setPluginNamespace( void SpecialSlicePluginDynamicCreator::setPluginNamespace(
const char* lib_namespace) { const char* lib_namespace) TRT_NOEXCEPT {
plugin_namespace_ = lib_namespace; plugin_namespace_ = lib_namespace;
} }
const char* SpecialSlicePluginDynamicCreator::getPluginNamespace() const { const char* SpecialSlicePluginDynamicCreator::getPluginNamespace() const
TRT_NOEXCEPT {
return plugin_namespace_.c_str(); return plugin_namespace_.c_str();
} }
......
paddle/fluid/inference/tensorrt/plugin/special_slice_plugin.h
浏览文件 @ b5aab4f0
...@@ -31,37 +31,38 @@ class SpecialSlicePluginDynamic : public DynamicPluginTensorRT { ...@@ -31,37 +31,38 @@ class SpecialSlicePluginDynamic : public DynamicPluginTensorRT {
SpecialSlicePluginDynamic(); SpecialSlicePluginDynamic();
SpecialSlicePluginDynamic(void const* serial_data, size_t serial_length); SpecialSlicePluginDynamic(void const* serial_data, size_t serial_length);
~SpecialSlicePluginDynamic(); ~SpecialSlicePluginDynamic();
nvinfer1::IPluginV2DynamicExt* clone() const override; nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override;
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs, int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs,
nvinfer1::IExprBuilder& exprBuilder) override; nvinfer1::IExprBuilder& exprBuilder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* inOut, const nvinfer1::PluginTensorDesc* inOut,
int nbInputs, int nbOutputs) override; int nbInputs,
int nbOutputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nbInputs, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nbOutputs) override; int nbOutputs) TRT_NOEXCEPT override;
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nbInputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const override; int nbOutputs) const TRT_NOEXCEPT override;
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* inputTypes, int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const override; int nbInputs) const TRT_NOEXCEPT override;
const char* getPluginType() const override; const char* getPluginType() const TRT_NOEXCEPT override;
int getNbOutputs() const override; int getNbOutputs() const TRT_NOEXCEPT override;
int initialize() override; int initialize() TRT_NOEXCEPT override;
void terminate() override; void terminate() TRT_NOEXCEPT override;
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
void destroy() override; void destroy() TRT_NOEXCEPT override;
private: private:
int axis_; int axis_;
...@@ -71,16 +72,17 @@ class SpecialSlicePluginDynamic : public DynamicPluginTensorRT { ...@@ -71,16 +72,17 @@ class SpecialSlicePluginDynamic : public DynamicPluginTensorRT {
class SpecialSlicePluginDynamicCreator : public nvinfer1::IPluginCreator { class SpecialSlicePluginDynamicCreator : public nvinfer1::IPluginCreator {
public: public:
SpecialSlicePluginDynamicCreator(); SpecialSlicePluginDynamicCreator();
const char* getPluginName() const override; const char* getPluginName() const TRT_NOEXCEPT override;
const char* getPluginVersion() const override; const char* getPluginVersion() const TRT_NOEXCEPT override;
const nvinfer1::PluginFieldCollection* getFieldNames() override; const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override;
nvinfer1::IPluginV2* createPlugin( nvinfer1::IPluginV2* createPlugin(const char* name,
const char* name, const nvinfer1::PluginFieldCollection* fc) override; const nvinfer1::PluginFieldCollection* fc)
nvinfer1::IPluginV2* deserializePlugin(const char* name, TRT_NOEXCEPT override;
const void* serial_data, nvinfer1::IPluginV2* deserializePlugin(
size_t serial_length) override; const char* name, const void* serial_data,
void setPluginNamespace(const char* lib_namespace) override; size_t serial_length) TRT_NOEXCEPT override;
const char* getPluginNamespace() const override; void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override;
const char* getPluginNamespace() const TRT_NOEXCEPT override;
private: private:
std::string plugin_namespace_; std::string plugin_namespace_;
......
paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -38,7 +38,7 @@ __device__ int upper_bound(T const* vals, int n, T const& key) { ...@@ -38,7 +38,7 @@ __device__ int upper_bound(T const* vals, int n, T const& key) {
} }
nvinfer1::Dims SplitPlugin::getOutputDimensions( nvinfer1::Dims SplitPlugin::getOutputDimensions(
int index, const nvinfer1::Dims* input_dims, int num_inputs) { int index, const nvinfer1::Dims* input_dims, int num_inputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(num_inputs, 1, PADDLE_ENFORCE_EQ(num_inputs, 1,
platform::errors::InvalidArgument( platform::errors::InvalidArgument(
"Invalid number of inputs of split TRT plugin. " "Invalid number of inputs of split TRT plugin. "
...@@ -66,7 +66,7 @@ void SplitPlugin::shareData(const SplitPlugin* another) { ...@@ -66,7 +66,7 @@ void SplitPlugin::shareData(const SplitPlugin* another) {
d_output_ptrs_.resize(another->d_output_ptrs_.size(), nullptr); d_output_ptrs_.resize(another->d_output_ptrs_.size(), nullptr);
} }
int SplitPlugin::initialize() { int SplitPlugin::initialize() TRT_NOEXCEPT {
PADDLE_ENFORCE_LE(axis_, nvinfer1::Dims::MAX_DIMS, PADDLE_ENFORCE_LE(axis_, nvinfer1::Dims::MAX_DIMS,
platform::errors::InvalidArgument( platform::errors::InvalidArgument(
"Axis dimension exceeds max dimension in TensorRT. " "Axis dimension exceeds max dimension in TensorRT. "
...@@ -98,7 +98,7 @@ int SplitPlugin::initialize() { ...@@ -98,7 +98,7 @@ int SplitPlugin::initialize() {
} }
// nothing to release according to initialize // nothing to release according to initialize
void SplitPlugin::terminate() {} void SplitPlugin::terminate() TRT_NOEXCEPT {}
// The following part of the code refers to onnx-tensorrt // The following part of the code refers to onnx-tensorrt
// https://github.com/onnx/onnx-tensorrt/blob/master/Split.cu // https://github.com/onnx/onnx-tensorrt/blob/master/Split.cu
...@@ -129,7 +129,7 @@ int SplitPlugin::enqueue(int batchSize, const void* const* inputs, ...@@ -129,7 +129,7 @@ int SplitPlugin::enqueue(int batchSize, const void* const* inputs,
void** outputs, void* workspace, cudaStream_t stream) { void** outputs, void* workspace, cudaStream_t stream) {
#else #else
void* const* outputs, void* workspace, void* const* outputs, void* workspace,
cudaStream_t stream) { cudaStream_t stream) TRT_NOEXCEPT {
#endif #endif
const int* d_segment_offsets_ptr = const int* d_segment_offsets_ptr =
thrust::raw_pointer_cast(&d_segment_offsets_[0]); thrust::raw_pointer_cast(&d_segment_offsets_[0]);
...@@ -155,14 +155,14 @@ int SplitPlugin::enqueue(int batchSize, const void* const* inputs, ...@@ -155,14 +155,14 @@ int SplitPlugin::enqueue(int batchSize, const void* const* inputs,
// Dynamic Plugin below. // Dynamic Plugin below.
#if IS_TRT_VERSION_GE(6000) #if IS_TRT_VERSION_GE(6000)
int SplitPluginDynamic::initialize() { return 0; } int SplitPluginDynamic::initialize() TRT_NOEXCEPT { return 0; }
size_t SplitPluginDynamic::getSerializationSize() const { size_t SplitPluginDynamic::getSerializationSize() const TRT_NOEXCEPT {
return SerializedSize(axis_) + SerializedSize(output_length_) + return SerializedSize(axis_) + SerializedSize(output_length_) +
SerializedSize(with_fp16_); SerializedSize(with_fp16_);
} }
void SplitPluginDynamic::serialize(void* buffer) const { void SplitPluginDynamic::serialize(void* buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, axis_); SerializeValue(&buffer, axis_);
SerializeValue(&buffer, output_length_); SerializeValue(&buffer, output_length_);
SerializeValue(&buffer, with_fp16_); SerializeValue(&buffer, with_fp16_);
...@@ -170,7 +170,7 @@ void SplitPluginDynamic::serialize(void* buffer) const { ...@@ -170,7 +170,7 @@ void SplitPluginDynamic::serialize(void* buffer) const {
nvinfer1::DimsExprs SplitPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs SplitPluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) { nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(nb_inputs, 1, PADDLE_ENFORCE_EQ(nb_inputs, 1,
platform::errors::InvalidArgument( platform::errors::InvalidArgument(
"The Split plugin should be only one input.")); "The Split plugin should be only one input."));
...@@ -188,7 +188,7 @@ nvinfer1::DimsExprs SplitPluginDynamic::getOutputDimensions( ...@@ -188,7 +188,7 @@ nvinfer1::DimsExprs SplitPluginDynamic::getOutputDimensions(
bool SplitPluginDynamic::supportsFormatCombination( bool SplitPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc* in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc* in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of split plugin should not be nullptr.")); "The input of split plugin should not be nullptr."));
...@@ -217,14 +217,16 @@ bool SplitPluginDynamic::supportsFormatCombination( ...@@ -217,14 +217,16 @@ bool SplitPluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType SplitPluginDynamic::getOutputDataType( nvinfer1::DataType SplitPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType* input_types, int nb_inputs) const { int index, const nvinfer1::DataType* input_types,
int nb_inputs) const TRT_NOEXCEPT {
return input_types[0]; return input_types[0];
} }
int SplitPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc* input_desc, int SplitPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc* input_desc,
const nvinfer1::PluginTensorDesc* output_desc, const nvinfer1::PluginTensorDesc* output_desc,
const void* const* inputs, void* const* outputs, const void* const* inputs, void* const* outputs,
void* workspace, cudaStream_t stream) { void* workspace,
cudaStream_t stream) TRT_NOEXCEPT {
auto input_dims = input_desc[0].dims; auto input_dims = input_desc[0].dims;
int outer_rows = 1; int outer_rows = 1;
int inner_cols = 1; int inner_cols = 1;
......
paddle/fluid/inference/tensorrt/plugin/split_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -39,43 +39,47 @@ class SplitPlugin : public PluginTensorRTV2Ext { ...@@ -39,43 +39,47 @@ class SplitPlugin : public PluginTensorRTV2Ext {
DeserializeValue(&serial_data, &serial_length, &output_length_); DeserializeValue(&serial_data, &serial_length, &output_length_);
} }
nvinfer1::IPluginV2Ext* clone() const override { nvinfer1::IPluginV2Ext* clone() const TRT_NOEXCEPT override {
SplitPlugin* ptr = new SplitPlugin(axis_, output_length_, with_fp16_); SplitPlugin* ptr = new SplitPlugin(axis_, output_length_, with_fp16_);
ptr->setPluginNamespace(this->getPluginNamespace()); ptr->setPluginNamespace(this->getPluginNamespace());
ptr->shareData(this); ptr->shareData(this);
return ptr; return ptr;
} }
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* input_types, int index, const nvinfer1::DataType* input_types,
int nb_inputs) const override { int nb_inputs) const TRT_NOEXCEPT override {
return input_types[0]; return input_types[0];
} }
const char* getPluginType() const override { return "split_plugin_v2ext"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return output_length_.size(); } return "split_plugin_v2ext";
}
int getNbOutputs() const TRT_NOEXCEPT override {
return output_length_.size();
}
nvinfer1::Dims getOutputDimensions(int index, nvinfer1::Dims getOutputDimensions(int index,
const nvinfer1::Dims* input_dims, const nvinfer1::Dims* input_dims,
int num_inputs) override; int num_inputs) TRT_NOEXCEPT override;
int initialize() override; int initialize() TRT_NOEXCEPT override;
void terminate() override; void terminate() TRT_NOEXCEPT override;
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
int enqueue(int batch_size, const void* const* inputs, void** outputs, int enqueue(int batch_size, const void* const* inputs, void** outputs,
#else #else
int enqueue(int batch_size, const void* const* inputs, void* const* outputs, int enqueue(int batch_size, const void* const* inputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream) override; void* workspace, cudaStream_t stream) TRT_NOEXCEPT override;
void destroy() override { delete this; } void destroy() TRT_NOEXCEPT override { delete this; }
protected: protected:
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
return SerializedSize(axis_) + SerializedSize(output_length_) + return SerializedSize(axis_) + SerializedSize(output_length_) +
getBaseSerializationSize(); getBaseSerializationSize();
} }
void serialize(void* buffer) const override { void serialize(void* buffer) const TRT_NOEXCEPT override {
serializeBase(buffer); serializeBase(buffer);
SerializeValue(&buffer, axis_); SerializeValue(&buffer, axis_);
SerializeValue(&buffer, output_length_); SerializeValue(&buffer, output_length_);
...@@ -98,32 +102,35 @@ class SplitPlugin : public PluginTensorRTV2Ext { ...@@ -98,32 +102,35 @@ class SplitPlugin : public PluginTensorRTV2Ext {
class SplitPluginCreator : public nvinfer1::IPluginCreator { class SplitPluginCreator : public nvinfer1::IPluginCreator {
public: public:
SplitPluginCreator() {} SplitPluginCreator() {}
const char* getPluginName() const override { return "split_plugin_v2ext"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "split_plugin_v2ext";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
const nvinfer1::PluginFieldCollection* getFieldNames() override { const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2* createPlugin( nvinfer1::IPluginV2* createPlugin(const char* name,
const char* name, const nvinfer1::PluginFieldCollection* fc) override { const nvinfer1::PluginFieldCollection* fc)
TRT_NOEXCEPT override {
// not implemented // not implemented
return nullptr; return nullptr;
} }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
auto plugin = new SplitPlugin(serial_data, serial_length); auto plugin = new SplitPlugin(serial_data, serial_length);
return plugin; return plugin;
} }
void setPluginNamespace(const char* lib_namespace) override { void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override {
plugin_namespace_ = lib_namespace; plugin_namespace_ = lib_namespace;
} }
const char* getPluginNamespace() const override { const char* getPluginNamespace() const TRT_NOEXCEPT override {
return plugin_namespace_.c_str(); return plugin_namespace_.c_str();
} }
...@@ -151,46 +158,51 @@ class SplitPluginDynamic : public DynamicPluginTensorRT { ...@@ -151,46 +158,51 @@ class SplitPluginDynamic : public DynamicPluginTensorRT {
DeserializeValue(&serial_data, &serial_length, &with_fp16_); DeserializeValue(&serial_data, &serial_length, &with_fp16_);
} }
nvinfer1::IPluginV2DynamicExt* clone() const override { nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override {
return new SplitPluginDynamic(axis_, output_length_, with_fp16_); return new SplitPluginDynamic(axis_, output_length_, with_fp16_);
} }
const char* getPluginType() const override { return "split_plugin"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return output_length_.size(); } return "split_plugin";
int initialize() override; }
int getNbOutputs() const TRT_NOEXCEPT override {
return output_length_.size();
}
int initialize() TRT_NOEXCEPT override;
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs, int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs,
nvinfer1::IExprBuilder& exprBuilder) override; nvinfer1::IExprBuilder& exprBuilder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* inOut, const nvinfer1::PluginTensorDesc* inOut,
int nbInputs, int nbOutputs) override; int nbInputs,
int nbOutputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nbInputs, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nbOutputs) override {} int nbOutputs) TRT_NOEXCEPT override {}
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nbInputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const override { int nbOutputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* inputTypes, int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const override; int nbInputs) const TRT_NOEXCEPT override;
void destroy() override { delete this; } void destroy() TRT_NOEXCEPT override { delete this; }
private: private:
int axis_; int axis_;
...@@ -200,31 +212,34 @@ class SplitPluginDynamic : public DynamicPluginTensorRT { ...@@ -200,31 +212,34 @@ class SplitPluginDynamic : public DynamicPluginTensorRT {
class SplitPluginDynamicCreator : public nvinfer1::IPluginCreator { class SplitPluginDynamicCreator : public nvinfer1::IPluginCreator {
public: public:
SplitPluginDynamicCreator() {} SplitPluginDynamicCreator() {}
const char* getPluginName() const override { return "split_plugin"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "split_plugin";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
const nvinfer1::PluginFieldCollection* getFieldNames() override { const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2* createPlugin( nvinfer1::IPluginV2* createPlugin(const char* name,
const char* name, const nvinfer1::PluginFieldCollection* fc) override { const nvinfer1::PluginFieldCollection* fc)
TRT_NOEXCEPT override {
return nullptr; return nullptr;
} }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
auto plugin = new SplitPluginDynamic(serial_data, serial_length); auto plugin = new SplitPluginDynamic(serial_data, serial_length);
return plugin; return plugin;
} }
void setPluginNamespace(const char* lib_namespace) override { void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override {
plugin_namespace_ = lib_namespace; plugin_namespace_ = lib_namespace;
} }
const char* getPluginNamespace() const override { const char* getPluginNamespace() const TRT_NOEXCEPT override {
return plugin_namespace_.c_str(); return plugin_namespace_.c_str();
} }
......
paddle/fluid/inference/tensorrt/plugin/stack_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -37,17 +37,19 @@ StackPluginDynamic::StackPluginDynamic(void const* serial_data, ...@@ -37,17 +37,19 @@ StackPluginDynamic::StackPluginDynamic(void const* serial_data,
StackPluginDynamic::~StackPluginDynamic() {} StackPluginDynamic::~StackPluginDynamic() {}
nvinfer1::IPluginV2DynamicExt* StackPluginDynamic::clone() const { nvinfer1::IPluginV2DynamicExt* StackPluginDynamic::clone() const TRT_NOEXCEPT {
return new StackPluginDynamic(axis_, num_stack_, with_fp16_); return new StackPluginDynamic(axis_, num_stack_, with_fp16_);
} }
const char* StackPluginDynamic::getPluginType() const { return "stack_plugin"; } const char* StackPluginDynamic::getPluginType() const TRT_NOEXCEPT {
return "stack_plugin";
}
int StackPluginDynamic::getNbOutputs() const { return 1; } int StackPluginDynamic::getNbOutputs() const TRT_NOEXCEPT { return 1; }
int StackPluginDynamic::initialize() { return 0; } int StackPluginDynamic::initialize() TRT_NOEXCEPT { return 0; }
size_t StackPluginDynamic::getSerializationSize() const { size_t StackPluginDynamic::getSerializationSize() const TRT_NOEXCEPT {
size_t serialize_size = 0; size_t serialize_size = 0;
serialize_size += SerializedSize(axis_); serialize_size += SerializedSize(axis_);
serialize_size += SerializedSize(num_stack_); serialize_size += SerializedSize(num_stack_);
...@@ -55,7 +57,7 @@ size_t StackPluginDynamic::getSerializationSize() const { ...@@ -55,7 +57,7 @@ size_t StackPluginDynamic::getSerializationSize() const {
return serialize_size; return serialize_size;
} }
void StackPluginDynamic::serialize(void* buffer) const { void StackPluginDynamic::serialize(void* buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, axis_); SerializeValue(&buffer, axis_);
SerializeValue(&buffer, num_stack_); SerializeValue(&buffer, num_stack_);
SerializeValue(&buffer, with_fp16_); SerializeValue(&buffer, with_fp16_);
...@@ -63,7 +65,7 @@ void StackPluginDynamic::serialize(void* buffer) const { ...@@ -63,7 +65,7 @@ void StackPluginDynamic::serialize(void* buffer) const {
nvinfer1::DimsExprs StackPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs StackPluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) { nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT {
nvinfer1::DimsExprs output(inputs[0]); nvinfer1::DimsExprs output(inputs[0]);
output.nbDims = inputs[0].nbDims + 1; output.nbDims = inputs[0].nbDims + 1;
...@@ -76,21 +78,22 @@ nvinfer1::DimsExprs StackPluginDynamic::getOutputDimensions( ...@@ -76,21 +78,22 @@ nvinfer1::DimsExprs StackPluginDynamic::getOutputDimensions(
void StackPluginDynamic::configurePlugin( void StackPluginDynamic::configurePlugin(
const nvinfer1::DynamicPluginTensorDesc* in, int nbInputs, const nvinfer1::DynamicPluginTensorDesc* in, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, int nbOutputs) {} const nvinfer1::DynamicPluginTensorDesc* out, int nbOutputs) TRT_NOEXCEPT {}
size_t StackPluginDynamic::getWorkspaceSize( size_t StackPluginDynamic::getWorkspaceSize(
const nvinfer1::PluginTensorDesc* inputs, int nbInputs, const nvinfer1::PluginTensorDesc* inputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, int nbOutputs) const { const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const TRT_NOEXCEPT {
return num_stack_ * sizeof(uintptr_t); return num_stack_ * sizeof(uintptr_t);
} }
void StackPluginDynamic::destroy() { delete this; } void StackPluginDynamic::destroy() TRT_NOEXCEPT { delete this; }
void StackPluginDynamic::terminate() {} void StackPluginDynamic::terminate() TRT_NOEXCEPT {}
bool StackPluginDynamic::supportsFormatCombination( bool StackPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc* in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc* in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of stack plugin should not be nullptr.")); "The input of stack plugin should not be nullptr."));
...@@ -118,7 +121,8 @@ bool StackPluginDynamic::supportsFormatCombination( ...@@ -118,7 +121,8 @@ bool StackPluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType StackPluginDynamic::getOutputDataType( nvinfer1::DataType StackPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType* input_types, int nb_inputs) const { int index, const nvinfer1::DataType* input_types,
int nb_inputs) const TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument( PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument(
"The index should be equal to 0")); "The index should be equal to 0"));
return input_types[0]; return input_types[0];
...@@ -139,7 +143,8 @@ __global__ void StackKernel(const T* const* input, T* output, int num_stack, ...@@ -139,7 +143,8 @@ __global__ void StackKernel(const T* const* input, T* output, int num_stack,
int StackPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc* input_desc, int StackPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc* input_desc,
const nvinfer1::PluginTensorDesc* output_desc, const nvinfer1::PluginTensorDesc* output_desc,
const void* const* inputs, void* const* outputs, const void* const* inputs, void* const* outputs,
void* workspace, cudaStream_t stream) { void* workspace,
cudaStream_t stream) TRT_NOEXCEPT {
auto input_dims = input_desc[0].dims; // (batch, seq, seq) auto input_dims = input_desc[0].dims; // (batch, seq, seq)
auto out_dims = output_desc[0].dims; // (batch, num_head, seq, seq) auto out_dims = output_desc[0].dims; // (batch, num_head, seq, seq)
auto out_num_dims = out_dims.nbDims; auto out_num_dims = out_dims.nbDims;
...@@ -195,19 +200,21 @@ int StackPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc* input_desc, ...@@ -195,19 +200,21 @@ int StackPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc* input_desc,
StackPluginDynamicCreator::StackPluginDynamicCreator() {} StackPluginDynamicCreator::StackPluginDynamicCreator() {}
const char* StackPluginDynamicCreator::getPluginName() const { const char* StackPluginDynamicCreator::getPluginName() const TRT_NOEXCEPT {
return "stack_plugin"; return "stack_plugin";
} }
const char* StackPluginDynamicCreator::getPluginVersion() const { return "1"; } const char* StackPluginDynamicCreator::getPluginVersion() const TRT_NOEXCEPT {
return "1";
}
const nvinfer1::PluginFieldCollection* const nvinfer1::PluginFieldCollection*
StackPluginDynamicCreator::getFieldNames() { StackPluginDynamicCreator::getFieldNames() TRT_NOEXCEPT {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2* StackPluginDynamicCreator::createPlugin( nvinfer1::IPluginV2* StackPluginDynamicCreator::createPlugin(
const char* name, const nvinfer1::PluginFieldCollection* fc) { const char* name, const nvinfer1::PluginFieldCollection* fc) TRT_NOEXCEPT {
int axis = -1; int axis = -1;
int num_stack = -1; int num_stack = -1;
bool with_fp16 = false; bool with_fp16 = false;
...@@ -230,16 +237,18 @@ nvinfer1::IPluginV2* StackPluginDynamicCreator::createPlugin( ...@@ -230,16 +237,18 @@ nvinfer1::IPluginV2* StackPluginDynamicCreator::createPlugin(
} }
nvinfer1::IPluginV2* StackPluginDynamicCreator::deserializePlugin( nvinfer1::IPluginV2* StackPluginDynamicCreator::deserializePlugin(
const char* name, const void* serial_data, size_t serial_length) { const char* name, const void* serial_data,
size_t serial_length) TRT_NOEXCEPT {
auto plugin = new StackPluginDynamic(serial_data, serial_length); auto plugin = new StackPluginDynamic(serial_data, serial_length);
return plugin; return plugin;
} }
void StackPluginDynamicCreator::setPluginNamespace(const char* lib_namespace) { void StackPluginDynamicCreator::setPluginNamespace(const char* lib_namespace)
TRT_NOEXCEPT {
plugin_namespace_ = lib_namespace; plugin_namespace_ = lib_namespace;
} }
const char* StackPluginDynamicCreator::getPluginNamespace() const { const char* StackPluginDynamicCreator::getPluginNamespace() const TRT_NOEXCEPT {
return plugin_namespace_.c_str(); return plugin_namespace_.c_str();
} }
......
paddle/fluid/inference/tensorrt/plugin/stack_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -31,37 +31,36 @@ class StackPluginDynamic : public DynamicPluginTensorRT { ...@@ -31,37 +31,36 @@ class StackPluginDynamic : public DynamicPluginTensorRT {
explicit StackPluginDynamic(int axis, int num_stack, bool with_fp16); explicit StackPluginDynamic(int axis, int num_stack, bool with_fp16);
StackPluginDynamic(void const* serial_data, size_t serial_length); StackPluginDynamic(void const* serial_data, size_t serial_length);
~StackPluginDynamic(); ~StackPluginDynamic();
nvinfer1::IPluginV2DynamicExt* clone() const override; nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override;
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs, int outputIndex, const nvinfer1::DimsExprs* inputs, int nbInputs,
nvinfer1::IExprBuilder& exprBuilder) override; nvinfer1::IExprBuilder& exprBuilder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* inOut, const nvinfer1::PluginTensorDesc* inOut,
int nbInputs, int nbOutputs) override; int nbInputs,
int nbOutputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nbInputs, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nbOutputs) override; int nbOutputs) TRT_NOEXCEPT override;
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nbInputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const override; int nbOutputs) const TRT_NOEXCEPT override;
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(
nvinfer1::DataType getOutputDataType(int index, int index, const nvinfer1::DataType* inputTypes,
const nvinfer1::DataType* inputTypes, int nbInputs) const TRT_NOEXCEPT override;
int nbInputs) const override; const char* getPluginType() const TRT_NOEXCEPT override;
int getNbOutputs() const TRT_NOEXCEPT override;
const char* getPluginType() const override; int initialize() TRT_NOEXCEPT override;
int getNbOutputs() const override; void terminate() TRT_NOEXCEPT override;
int initialize() override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void terminate() override; void serialize(void* buffer) const TRT_NOEXCEPT override;
size_t getSerializationSize() const override; void destroy() TRT_NOEXCEPT override;
void serialize(void* buffer) const override;
void destroy() override;
private: private:
int axis_; int axis_;
...@@ -71,16 +70,17 @@ class StackPluginDynamic : public DynamicPluginTensorRT { ...@@ -71,16 +70,17 @@ class StackPluginDynamic : public DynamicPluginTensorRT {
class StackPluginDynamicCreator : public nvinfer1::IPluginCreator { class StackPluginDynamicCreator : public nvinfer1::IPluginCreator {
public: public:
StackPluginDynamicCreator(); StackPluginDynamicCreator();
const char* getPluginName() const override; const char* getPluginName() const TRT_NOEXCEPT override;
const char* getPluginVersion() const override; const char* getPluginVersion() const TRT_NOEXCEPT override;
const nvinfer1::PluginFieldCollection* getFieldNames() override; const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override;
nvinfer1::IPluginV2* createPlugin( nvinfer1::IPluginV2* createPlugin(const char* name,
const char* name, const nvinfer1::PluginFieldCollection* fc) override; const nvinfer1::PluginFieldCollection* fc)
nvinfer1::IPluginV2* deserializePlugin(const char* name, TRT_NOEXCEPT override;
const void* serial_data, nvinfer1::IPluginV2* deserializePlugin(
size_t serial_length) override; const char* name, const void* serial_data,
void setPluginNamespace(const char* lib_namespace) override; size_t serial_length) TRT_NOEXCEPT override;
const char* getPluginNamespace() const override; void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override;
const char* getPluginNamespace() const TRT_NOEXCEPT override;
private: private:
std::string plugin_namespace_; std::string plugin_namespace_;
......
paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -23,11 +23,11 @@ namespace inference { ...@@ -23,11 +23,11 @@ namespace inference {
namespace tensorrt { namespace tensorrt {
namespace plugin { namespace plugin {
int SwishPlugin::initialize() { return 0; } int SwishPlugin::initialize() TRT_NOEXCEPT { return 0; }
nvinfer1::Dims SwishPlugin::getOutputDimensions(int index, nvinfer1::Dims SwishPlugin::getOutputDimensions(int index,
const nvinfer1::Dims *inputDims, const nvinfer1::Dims *inputDims,
int nbInputs) { int nbInputs) TRT_NOEXCEPT {
assert(nbInputs == 1); assert(nbInputs == 1);
assert(index < this->getNbOutputs()); assert(index < this->getNbOutputs());
nvinfer1::Dims const &input_dims = inputDims[0]; nvinfer1::Dims const &input_dims = inputDims[0];
...@@ -83,12 +83,12 @@ int SwishPlugin::enqueue(int batch_size, const void *const *inputs, ...@@ -83,12 +83,12 @@ int SwishPlugin::enqueue(int batch_size, const void *const *inputs,
void **outputs, void *workspace, cudaStream_t stream) { void **outputs, void *workspace, cudaStream_t stream) {
#else #else
void *const *outputs, void *workspace, void *const *outputs, void *workspace,
cudaStream_t stream) { cudaStream_t stream) TRT_NOEXCEPT {
#endif #endif
// input dims is CHW. // input dims is CHW.
const auto &input_dims = this->getInputDims(0); const auto &input_dims = this->getInputDims(0);
const float *input = reinterpret_cast<const float *>(inputs[0]); const float *input = reinterpret_cast<const float *>(inputs[0]);
float *output = reinterpret_cast<float **>(outputs)[0]; float *output = reinterpret_cast<float *const *>(outputs)[0];
int num = batch_size; int num = batch_size;
for (int i = 0; i < input_dims.nbDims; i++) { for (int i = 0; i < input_dims.nbDims; i++) {
num *= input_dims.d[i]; num *= input_dims.d[i];
...@@ -103,29 +103,29 @@ int SwishPlugin::enqueue(int batch_size, const void *const *inputs, ...@@ -103,29 +103,29 @@ int SwishPlugin::enqueue(int batch_size, const void *const *inputs,
// Dynamic Plugin below. // Dynamic Plugin below.
#if IS_TRT_VERSION_GE(6000) #if IS_TRT_VERSION_GE(6000)
int SwishPluginDynamic::initialize() { int SwishPluginDynamic::initialize() TRT_NOEXCEPT {
getPluginNamespace(); getPluginNamespace();
return 0; return 0;
} }
size_t SwishPluginDynamic::getSerializationSize() const { size_t SwishPluginDynamic::getSerializationSize() const TRT_NOEXCEPT {
return SerializedSize(beta_) + SerializedSize(with_fp16_); return SerializedSize(beta_) + SerializedSize(with_fp16_);
} }
void SwishPluginDynamic::serialize(void *buffer) const { void SwishPluginDynamic::serialize(void *buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, beta_); SerializeValue(&buffer, beta_);
SerializeValue(&buffer, with_fp16_); SerializeValue(&buffer, with_fp16_);
} }
nvinfer1::DimsExprs SwishPluginDynamic::getOutputDimensions( nvinfer1::DimsExprs SwishPluginDynamic::getOutputDimensions(
int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs *inputs, int nb_inputs,
nvinfer1::IExprBuilder &expr_builder) { nvinfer1::IExprBuilder &expr_builder) TRT_NOEXCEPT {
return inputs[0]; return inputs[0];
} }
bool SwishPluginDynamic::supportsFormatCombination( bool SwishPluginDynamic::supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc *in_out, int nb_inputs,
int nb_outputs) { int nb_outputs) TRT_NOEXCEPT {
PADDLE_ENFORCE_NOT_NULL( PADDLE_ENFORCE_NOT_NULL(
in_out, platform::errors::InvalidArgument( in_out, platform::errors::InvalidArgument(
"The input of swish plugin shoule not be nullptr.")); "The input of swish plugin shoule not be nullptr."));
...@@ -154,7 +154,8 @@ bool SwishPluginDynamic::supportsFormatCombination( ...@@ -154,7 +154,8 @@ bool SwishPluginDynamic::supportsFormatCombination(
} }
nvinfer1::DataType SwishPluginDynamic::getOutputDataType( nvinfer1::DataType SwishPluginDynamic::getOutputDataType(
int index, const nvinfer1::DataType *input_types, int nb_inputs) const { int index, const nvinfer1::DataType *input_types,
int nb_inputs) const TRT_NOEXCEPT {
PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument( PADDLE_ENFORCE_EQ(index, 0, platform::errors::InvalidArgument(
"The Swish Plugin only has one input, so the " "The Swish Plugin only has one input, so the "
"index value should be 0, but get %d.", "index value should be 0, but get %d.",
...@@ -165,7 +166,8 @@ nvinfer1::DataType SwishPluginDynamic::getOutputDataType( ...@@ -165,7 +166,8 @@ nvinfer1::DataType SwishPluginDynamic::getOutputDataType(
int SwishPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc *input_desc, int SwishPluginDynamic::enqueue(const nvinfer1::PluginTensorDesc *input_desc,
const nvinfer1::PluginTensorDesc *output_desc, const nvinfer1::PluginTensorDesc *output_desc,
const void *const *inputs, void *const *outputs, const void *const *inputs, void *const *outputs,
void *workspace, cudaStream_t stream) { void *workspace,
cudaStream_t stream) TRT_NOEXCEPT {
auto input_dims = input_desc[0].dims; auto input_dims = input_desc[0].dims;
size_t num = ProductDim(input_dims); size_t num = ProductDim(input_dims);
int threads = 1024; int threads = 1024;
......
paddle/fluid/inference/tensorrt/plugin/swish_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -31,11 +31,11 @@ class SwishPlugin : public PluginTensorRT { ...@@ -31,11 +31,11 @@ class SwishPlugin : public PluginTensorRT {
float beta_; float beta_;
public: public:
size_t getSerializationSize() const override { size_t getSerializationSize() const TRT_NOEXCEPT override {
return getBaseSerializationSize() + SerializedSize(beta_); return getBaseSerializationSize() + SerializedSize(beta_);
} }
void serialize(void* buffer) const override { void serialize(void* buffer) const TRT_NOEXCEPT override {
serializeBase(buffer); serializeBase(buffer);
SerializeValue(&buffer, beta_); SerializeValue(&buffer, beta_);
} }
...@@ -53,33 +53,37 @@ class SwishPlugin : public PluginTensorRT { ...@@ -53,33 +53,37 @@ class SwishPlugin : public PluginTensorRT {
~SwishPlugin() {} ~SwishPlugin() {}
int initialize() override; int initialize() TRT_NOEXCEPT override;
SwishPlugin* clone() const override { SwishPlugin* clone() const TRT_NOEXCEPT override {
return new SwishPlugin(beta_, with_fp16_); return new SwishPlugin(beta_, with_fp16_);
} }
const char* getPluginType() const override { return "swish_plugin"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "swish_plugin";
}
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs,
int nbInputDims) override; int nbInputDims) TRT_NOEXCEPT override;
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
int enqueue(int batchSize, const void* const* inputs, void** outputs, int enqueue(int batchSize, const void* const* inputs, void** outputs,
#else #else
int enqueue(int batchSize, const void* const* inputs, void* const* outputs, int enqueue(int batchSize, const void* const* inputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream) override; void* workspace, cudaStream_t stream) TRT_NOEXCEPT override;
}; };
class SwishPluginCreator : public TensorRTPluginCreator { class SwishPluginCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "swish_plugin"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "swish_plugin";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new SwishPlugin(serial_data, serial_length); return new SwishPlugin(serial_data, serial_length);
} }
}; };
...@@ -96,46 +100,49 @@ class SwishPluginDynamic : public DynamicPluginTensorRT { ...@@ -96,46 +100,49 @@ class SwishPluginDynamic : public DynamicPluginTensorRT {
DeserializeValue(&serialData, &serialLength, &beta_); DeserializeValue(&serialData, &serialLength, &beta_);
DeserializeValue(&serialData, &serialLength, &with_fp16_); DeserializeValue(&serialData, &serialLength, &with_fp16_);
} }
nvinfer1::IPluginV2DynamicExt* clone() const override { nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT override {
return new SwishPluginDynamic(beta_, with_fp16_); return new SwishPluginDynamic(beta_, with_fp16_);
} }
const char* getPluginType() const override { return "swish_plugin_dynamic"; } const char* getPluginType() const TRT_NOEXCEPT override {
int getNbOutputs() const override { return 1; } return "swish_plugin_dynamic";
int initialize() override; }
int getNbOutputs() const TRT_NOEXCEPT override { return 1; }
int initialize() TRT_NOEXCEPT override;
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
nvinfer1::DimsExprs getOutputDimensions( nvinfer1::DimsExprs getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) override; nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT override;
bool supportsFormatCombination(int pos, bool supportsFormatCombination(int pos,
const nvinfer1::PluginTensorDesc* inOut, const nvinfer1::PluginTensorDesc* inOut,
int nbInputs, int nbOutputs) override; int nbInputs,
int nbOutputs) TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nbInputs, int nbInputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nbOutputs) override {} int nbOutputs) TRT_NOEXCEPT override {}
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nbInputs, int nbInputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nbOutputs) const override { int nbOutputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, int enqueue(const nvinfer1::PluginTensorDesc* inputDesc,
const nvinfer1::PluginTensorDesc* outputDesc, const nvinfer1::PluginTensorDesc* outputDesc,
const void* const* inputs, void* const* outputs, void* workspace, const void* const* inputs, void* const* outputs, void* workspace,
cudaStream_t stream) override; cudaStream_t stream) TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* inputTypes, int index, const nvinfer1::DataType* inputTypes,
int nbInputs) const override; int nbInputs) const TRT_NOEXCEPT override;
void destroy() override { delete this; } void destroy() TRT_NOEXCEPT override { delete this; }
private: private:
float beta_; float beta_;
...@@ -143,13 +150,15 @@ class SwishPluginDynamic : public DynamicPluginTensorRT { ...@@ -143,13 +150,15 @@ class SwishPluginDynamic : public DynamicPluginTensorRT {
class SwishPluginDynamicCreator : public TensorRTPluginCreator { class SwishPluginDynamicCreator : public TensorRTPluginCreator {
public: public:
const char* getPluginName() const override { return "swish_plugin_dynamic"; } const char* getPluginName() const TRT_NOEXCEPT override {
return "swish_plugin_dynamic";
}
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
nvinfer1::IPluginV2* deserializePlugin(const char* name, nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) override { size_t serial_length) TRT_NOEXCEPT override {
return new SwishPluginDynamic(serial_data, serial_length); return new SwishPluginDynamic(serial_data, serial_length);
} }
}; };
......
paddle/fluid/inference/tensorrt/plugin/trt_plugin.cc
浏览文件 @ b5aab4f0
...@@ -60,8 +60,8 @@ size_t PluginTensorRT::getBaseSerializationSize() const { ...@@ -60,8 +60,8 @@ size_t PluginTensorRT::getBaseSerializationSize() const {
return SeriaSize(input_dims_, data_type_, data_format_, with_fp16_); return SeriaSize(input_dims_, data_type_, data_format_, with_fp16_);
} }
bool PluginTensorRT::supportsFormat(nvinfer1::DataType type, bool PluginTensorRT::supportsFormat(
nvinfer1::PluginFormat format) const { nvinfer1::DataType type, nvinfer1::PluginFormat format) const TRT_NOEXCEPT {
return ((type == nvinfer1::DataType::kFLOAT) && return ((type == nvinfer1::DataType::kFLOAT) &&
(format == nvinfer1::PluginFormat::kLINEAR)); (format == nvinfer1::PluginFormat::kLINEAR));
} }
...@@ -69,7 +69,7 @@ bool PluginTensorRT::supportsFormat(nvinfer1::DataType type, ...@@ -69,7 +69,7 @@ bool PluginTensorRT::supportsFormat(nvinfer1::DataType type,
void PluginTensorRT::configureWithFormat( void PluginTensorRT::configureWithFormat(
const nvinfer1::Dims* input_dims, int num_inputs, const nvinfer1::Dims* input_dims, int num_inputs,
const nvinfer1::Dims* output_dims, int num_outputs, nvinfer1::DataType type, const nvinfer1::Dims* output_dims, int num_outputs, nvinfer1::DataType type,
nvinfer1::PluginFormat format, int max_batch_size) { nvinfer1::PluginFormat format, int max_batch_size) TRT_NOEXCEPT {
data_type_ = type; data_type_ = type;
data_format_ = format; data_format_ = format;
input_dims_.assign(input_dims, input_dims + num_inputs); input_dims_.assign(input_dims, input_dims + num_inputs);
...@@ -95,26 +95,28 @@ void PluginTensorRTV2Ext::configurePlugin( ...@@ -95,26 +95,28 @@ void PluginTensorRTV2Ext::configurePlugin(
const nvinfer1::DataType* input_types, const nvinfer1::DataType* input_types,
const nvinfer1::DataType* output_types, const bool* input_is_broadcast, const nvinfer1::DataType* output_types, const bool* input_is_broadcast,
const bool* output_is_broadcast, nvinfer1::PluginFormat float_format, const bool* output_is_broadcast, nvinfer1::PluginFormat float_format,
int32_t max_batch_size) { int32_t max_batch_size) TRT_NOEXCEPT {
input_dims_.assign(input_dims, input_dims + nb_inputs); input_dims_.assign(input_dims, input_dims + nb_inputs);
data_format_ = float_format; data_format_ = float_format;
data_type_ = input_types[0]; data_type_ = input_types[0];
} }
const nvinfer1::PluginFieldCollection* TensorRTPluginCreator::getFieldNames() { const nvinfer1::PluginFieldCollection* TensorRTPluginCreator::getFieldNames()
TRT_NOEXCEPT {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2* TensorRTPluginCreator::createPlugin( nvinfer1::IPluginV2* TensorRTPluginCreator::createPlugin(
const char* name, const nvinfer1::PluginFieldCollection* fc) { const char* name, const nvinfer1::PluginFieldCollection* fc) TRT_NOEXCEPT {
return nullptr; return nullptr;
} }
void TensorRTPluginCreator::setPluginNamespace(const char* lib_namespace) { void TensorRTPluginCreator::setPluginNamespace(const char* lib_namespace)
TRT_NOEXCEPT {
plugin_namespace_ = lib_namespace; plugin_namespace_ = lib_namespace;
} }
const char* TensorRTPluginCreator::getPluginNamespace() const { const char* TensorRTPluginCreator::getPluginNamespace() const TRT_NOEXCEPT {
return plugin_namespace_.c_str(); return plugin_namespace_.c_str();
} }
......
paddle/fluid/inference/tensorrt/plugin/trt_plugin.h
浏览文件 @ b5aab4f0
...@@ -64,35 +64,35 @@ class PluginTensorRT : public nvinfer1::IPluginV2 { ...@@ -64,35 +64,35 @@ class PluginTensorRT : public nvinfer1::IPluginV2 {
nvinfer1::PluginFormat getDataFormat() const { return data_format_; } nvinfer1::PluginFormat getDataFormat() const { return data_format_; }
// IPluginV2 // IPluginV2
virtual const char* getPluginType() const = 0; virtual const char* getPluginType() const TRT_NOEXCEPT = 0;
virtual const char* getPluginVersion() const { return "1"; } virtual const char* getPluginVersion() const TRT_NOEXCEPT { return "1"; }
int getNbOutputs() const { return 1; } int getNbOutputs() const TRT_NOEXCEPT { return 1; }
virtual nvinfer1::Dims getOutputDimensions(int index, virtual nvinfer1::Dims getOutputDimensions(int index,
const nvinfer1::Dims* input_dims, const nvinfer1::Dims* input_dims,
int num_inputs) = 0; int num_inputs) TRT_NOEXCEPT = 0;
// Check format support. The default is FLOAT32 and kLINEAR. // Check format support. The default is FLOAT32 and kLINEAR.
bool supportsFormat(nvinfer1::DataType type, bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format)
nvinfer1::PluginFormat format) const override; const TRT_NOEXCEPT override;
// Configure the layer // Configure the layer
void configureWithFormat(const nvinfer1::Dims* input_dims, int num_inputs, void configureWithFormat(const nvinfer1::Dims* input_dims, int num_inputs,
const nvinfer1::Dims* output_dims, int num_outputs, const nvinfer1::Dims* output_dims, int num_outputs,
nvinfer1::DataType type, nvinfer1::DataType type,
nvinfer1::PluginFormat format, nvinfer1::PluginFormat format,
int max_batch_size) override; int max_batch_size) TRT_NOEXCEPT override;
// Initialize the layer for execution. // Initialize the layer for execution.
int initialize() override { return 0; } int initialize() TRT_NOEXCEPT override { return 0; }
// Shutdown the layer. This is called when the engine is destroyed // Shutdown the layer. This is called when the engine is destroyed
void terminate() override {} void terminate() TRT_NOEXCEPT override {}
// Find the workspace size required by the layer // Find the workspace size required by the layer
size_t getWorkspaceSize(int) const override { return 0; } size_t getWorkspaceSize(int) const TRT_NOEXCEPT override { return 0; }
// Execute the layer // Execute the layer
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
...@@ -101,25 +101,27 @@ class PluginTensorRT : public nvinfer1::IPluginV2 { ...@@ -101,25 +101,27 @@ class PluginTensorRT : public nvinfer1::IPluginV2 {
virtual int enqueue(int batch_size, const void* const* inputs, virtual int enqueue(int batch_size, const void* const* inputs,
void* const* outputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream) = 0; void* workspace, cudaStream_t stream) TRT_NOEXCEPT = 0;
// Find the size of the serialization buffer required // Find the size of the serialization buffer required
virtual size_t getSerializationSize() const = 0; virtual size_t getSerializationSize() const TRT_NOEXCEPT = 0;
// Serialize the layer config to buffer. // Serialize the layer config to buffer.
// TensorRT will call this func to serialize the configuration of TensorRT // TensorRT will call this func to serialize the configuration of TensorRT
// engine. It should not be called by users. // engine. It should not be called by users.
virtual void serialize(void* buffer) const = 0; virtual void serialize(void* buffer) const TRT_NOEXCEPT = 0;
void destroy() override { delete this; } void destroy() TRT_NOEXCEPT override { delete this; }
virtual nvinfer1::IPluginV2* clone() const = 0; virtual nvinfer1::IPluginV2* clone() const TRT_NOEXCEPT = 0;
void setPluginNamespace(const char* plugin_namespace) override { void setPluginNamespace(const char* plugin_namespace) TRT_NOEXCEPT override {
namespace_ = plugin_namespace; namespace_ = plugin_namespace;
} }
const char* getPluginNamespace() const override { return namespace_.c_str(); } const char* getPluginNamespace() const TRT_NOEXCEPT override {
return namespace_.c_str();
}
protected: protected:
// Deserialize input_dims, max_batch_size, data_type, data_format // Deserialize input_dims, max_batch_size, data_type, data_format
...@@ -155,15 +157,16 @@ class PluginTensorRTV2Ext : public nvinfer1::IPluginV2Ext { ...@@ -155,15 +157,16 @@ class PluginTensorRTV2Ext : public nvinfer1::IPluginV2Ext {
// The Func in IPluginV2Ext // The Func in IPluginV2Ext
virtual nvinfer1::DataType getOutputDataType( virtual nvinfer1::DataType getOutputDataType(
int index, const nvinfer1::DataType* input_types, int index, const nvinfer1::DataType* input_types,
int nb_inputs) const = 0; int nb_inputs) const TRT_NOEXCEPT = 0;
virtual bool isOutputBroadcastAcrossBatch(int32_t output_index, virtual bool isOutputBroadcastAcrossBatch(
const bool* input_is_broadcasted, int32_t output_index, const bool* input_is_broadcasted,
int32_t nb_inputs) const { int32_t nb_inputs) const TRT_NOEXCEPT {
return false; return false;
} }
virtual bool canBroadcastInputAcrossBatch(int32_t input_index) const { virtual bool canBroadcastInputAcrossBatch(int32_t input_index) const
TRT_NOEXCEPT {
return false; return false;
} }
...@@ -174,37 +177,37 @@ class PluginTensorRTV2Ext : public nvinfer1::IPluginV2Ext { ...@@ -174,37 +177,37 @@ class PluginTensorRTV2Ext : public nvinfer1::IPluginV2Ext {
const bool* input_is_broadcast, const bool* input_is_broadcast,
const bool* output_is_broadcast, const bool* output_is_broadcast,
nvinfer1::PluginFormat float_format, nvinfer1::PluginFormat float_format,
int32_t max_batch_size) override; int32_t max_batch_size) TRT_NOEXCEPT override;
virtual IPluginV2Ext* clone() const = 0; virtual IPluginV2Ext* clone() const TRT_NOEXCEPT = 0;
void attachToContext(cudnnContext*, cublasContext*, void attachToContext(cudnnContext*, cublasContext*,
nvinfer1::IGpuAllocator*) override {} nvinfer1::IGpuAllocator*) TRT_NOEXCEPT override {}
void detachFromContext() override {} void detachFromContext() TRT_NOEXCEPT override {}
// The Func in IPluginV2 // The Func in IPluginV2
virtual const char* getPluginType() const = 0; virtual const char* getPluginType() const TRT_NOEXCEPT = 0;
const char* getPluginVersion() const override { return "1"; } const char* getPluginVersion() const TRT_NOEXCEPT override { return "1"; }
virtual int32_t getNbOutputs() const { return 1; } virtual int32_t getNbOutputs() const TRT_NOEXCEPT { return 1; }
virtual nvinfer1::Dims getOutputDimensions(int32_t index, virtual nvinfer1::Dims getOutputDimensions(int32_t index,
const nvinfer1::Dims* inputs, const nvinfer1::Dims* inputs,
int32_t nb_input) = 0; int32_t nb_input) TRT_NOEXCEPT = 0;
// Check format support. The default is FLOAT32 and NCHW. // Check format support. The default is FLOAT32 and NCHW.
bool supportsFormat(nvinfer1::DataType type, bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format)
nvinfer1::PluginFormat format) const override { const TRT_NOEXCEPT override {
return ((type == nvinfer1::DataType::kFLOAT) && return ((type == nvinfer1::DataType::kFLOAT) &&
(format == nvinfer1::PluginFormat::kLINEAR)); (format == nvinfer1::PluginFormat::kLINEAR));
} }
// Initialize the layer for execution. // Initialize the layer for execution.
// This is called when the engine is created. // This is called when the engine is created.
int initialize() override { return 0; } int initialize() TRT_NOEXCEPT override { return 0; }
// Shutdown the layer. This is called when the engine is destroyed // Shutdown the layer. This is called when the engine is destroyed
void terminate() override {} void terminate() TRT_NOEXCEPT override {}
// Find the workspace size required by the layer // Find the workspace size required by the layer
size_t getWorkspaceSize(int) const override { return 0; } size_t getWorkspaceSize(int) const TRT_NOEXCEPT override { return 0; }
// Execute the layer // Execute the layer
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
...@@ -213,23 +216,23 @@ class PluginTensorRTV2Ext : public nvinfer1::IPluginV2Ext { ...@@ -213,23 +216,23 @@ class PluginTensorRTV2Ext : public nvinfer1::IPluginV2Ext {
virtual int enqueue(int batch_size, const void* const* inputs, virtual int enqueue(int batch_size, const void* const* inputs,
void* const* outputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream) = 0; void* workspace, cudaStream_t stream) TRT_NOEXCEPT = 0;
// Find the size of the serialization buffer required // Find the size of the serialization buffer required
virtual size_t getSerializationSize() const = 0; virtual size_t getSerializationSize() const TRT_NOEXCEPT = 0;
// Serialize the layer config to buffer. // Serialize the layer config to buffer.
// TensorRT will call this func to serialize the configuration of TensorRT // TensorRT will call this func to serialize the configuration of TensorRT
// engine. It should not be called by users. // engine. It should not be called by users.
virtual void serialize(void* buffer) const = 0; virtual void serialize(void* buffer) const TRT_NOEXCEPT = 0;
virtual void destroy() = 0; virtual void destroy() TRT_NOEXCEPT = 0;
void setPluginNamespace(const char* plugin_namespace) override { void setPluginNamespace(const char* plugin_namespace) TRT_NOEXCEPT override {
name_space_ = plugin_namespace; name_space_ = plugin_namespace;
} }
const char* getPluginNamespace() const override { const char* getPluginNamespace() const TRT_NOEXCEPT override {
return name_space_.c_str(); return name_space_.c_str();
} }
...@@ -256,52 +259,52 @@ class DynamicPluginTensorRT : public nvinfer1::IPluginV2DynamicExt { ...@@ -256,52 +259,52 @@ class DynamicPluginTensorRT : public nvinfer1::IPluginV2DynamicExt {
DynamicPluginTensorRT(const void* serialized_data, size_t length) {} DynamicPluginTensorRT(const void* serialized_data, size_t length) {}
// The Func in IPluginExt or IpluginExtV2 // The Func in IPluginExt or IpluginExtV2
virtual const char* getPluginVersion() const { return "1"; } virtual const char* getPluginVersion() const TRT_NOEXCEPT { return "1"; }
virtual const char* getPluginType() const = 0; virtual const char* getPluginType() const TRT_NOEXCEPT = 0;
int getNbOutputs() const { return 1; } int getNbOutputs() const TRT_NOEXCEPT { return 1; }
int initialize() override { return 0; } int initialize() TRT_NOEXCEPT override { return 0; }
void terminate() override{}; void terminate() TRT_NOEXCEPT override{};
virtual size_t getSerializationSize() const = 0; virtual size_t getSerializationSize() const TRT_NOEXCEPT = 0;
virtual void serialize(void* buffer) const = 0; virtual void serialize(void* buffer) const TRT_NOEXCEPT = 0;
// The Func in IPluginV2 // The Func in IPluginV2
nvinfer1::IPluginV2DynamicExt* clone() const = 0; nvinfer1::IPluginV2DynamicExt* clone() const TRT_NOEXCEPT = 0;
virtual nvinfer1::DimsExprs getOutputDimensions( virtual nvinfer1::DimsExprs getOutputDimensions(
int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs, int output_index, const nvinfer1::DimsExprs* inputs, int nb_inputs,
nvinfer1::IExprBuilder& expr_builder) = 0; // NOLINT nvinfer1::IExprBuilder& expr_builder) TRT_NOEXCEPT = 0; // NOLINT
virtual bool supportsFormatCombination( virtual bool supportsFormatCombination(
int pos, const nvinfer1::PluginTensorDesc* in_out, int nb_inputs, int pos, const nvinfer1::PluginTensorDesc* in_out, int nb_inputs,
int nb_outputs) = 0; int nb_outputs) TRT_NOEXCEPT = 0;
virtual void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in, virtual void configurePlugin(const nvinfer1::DynamicPluginTensorDesc* in,
int nb_inputs, int nb_inputs,
const nvinfer1::DynamicPluginTensorDesc* out, const nvinfer1::DynamicPluginTensorDesc* out,
int nb_outputs) = 0; int nb_outputs) TRT_NOEXCEPT = 0;
size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs, size_t getWorkspaceSize(const nvinfer1::PluginTensorDesc* inputs,
int nb_inputs, int nb_inputs,
const nvinfer1::PluginTensorDesc* outputs, const nvinfer1::PluginTensorDesc* outputs,
int nb_outputs) const override { int nb_outputs) const TRT_NOEXCEPT override {
return 0; return 0;
} }
virtual int enqueue(const nvinfer1::PluginTensorDesc* input_desc, virtual int enqueue(const nvinfer1::PluginTensorDesc* input_desc,
const nvinfer1::PluginTensorDesc* output_desc, const nvinfer1::PluginTensorDesc* output_desc,
const void* const* inputs, void* const* outputs, const void* const* inputs, void* const* outputs,
void* workspace, cudaStream_t stream) = 0; void* workspace, cudaStream_t stream) TRT_NOEXCEPT = 0;
virtual nvinfer1::DataType getOutputDataType( virtual nvinfer1::DataType getOutputDataType(
int index, const nvinfer1::DataType* input_types, int index, const nvinfer1::DataType* input_types,
int nb_inputs) const = 0; int nb_inputs) const TRT_NOEXCEPT = 0;
void setPluginNamespace(const char* plugin_namespace) override { void setPluginNamespace(const char* plugin_namespace) TRT_NOEXCEPT override {
name_space_ = plugin_namespace; name_space_ = plugin_namespace;
} }
const char* getPluginNamespace() const override { const char* getPluginNamespace() const TRT_NOEXCEPT override {
return name_space_.c_str(); return name_space_.c_str();
} }
virtual void destroy() = 0; virtual void destroy() TRT_NOEXCEPT = 0;
protected: protected:
void deserializeBase(void const*& serial_data, // NOLINT void deserializeBase(void const*& serial_data, // NOLINT
...@@ -320,22 +323,23 @@ class TensorRTPluginCreator : public nvinfer1::IPluginCreator { ...@@ -320,22 +323,23 @@ class TensorRTPluginCreator : public nvinfer1::IPluginCreator {
public: public:
TensorRTPluginCreator() = default; TensorRTPluginCreator() = default;
virtual const char* getPluginName() const = 0; virtual const char* getPluginName() const TRT_NOEXCEPT = 0;
virtual const char* getPluginVersion() const = 0; virtual const char* getPluginVersion() const TRT_NOEXCEPT = 0;
const nvinfer1::PluginFieldCollection* getFieldNames() override; const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override;
nvinfer1::IPluginV2* createPlugin( nvinfer1::IPluginV2* createPlugin(const char* name,
const char* name, const nvinfer1::PluginFieldCollection* fc) override; const nvinfer1::PluginFieldCollection* fc)
TRT_NOEXCEPT override;
virtual nvinfer1::IPluginV2* deserializePlugin(const char* name, virtual nvinfer1::IPluginV2* deserializePlugin(
const void* serial_data, const char* name, const void* serial_data,
size_t serial_length) = 0; size_t serial_length) TRT_NOEXCEPT = 0;
void setPluginNamespace(const char* lib_namespace) override; void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override;
const char* getPluginNamespace() const override; const char* getPluginNamespace() const TRT_NOEXCEPT override;
private: private:
std::string plugin_namespace_; std::string plugin_namespace_;
......
paddle/fluid/inference/tensorrt/plugin/yolo_box_op_plugin.cu
浏览文件 @ b5aab4f0
...@@ -70,15 +70,16 @@ YoloBoxPlugin::~YoloBoxPlugin() { ...@@ -70,15 +70,16 @@ YoloBoxPlugin::~YoloBoxPlugin() {
} }
} }
const char* YoloBoxPlugin::getPluginType() const { return "yolo_box_plugin"; } const char* YoloBoxPlugin::getPluginType() const TRT_NOEXCEPT {
return "yolo_box_plugin";
}
const char* YoloBoxPlugin::getPluginVersion() const { return "1"; } const char* YoloBoxPlugin::getPluginVersion() const TRT_NOEXCEPT { return "1"; }
int YoloBoxPlugin::getNbOutputs() const { return 2; } int YoloBoxPlugin::getNbOutputs() const TRT_NOEXCEPT { return 2; }
nvinfer1::Dims YoloBoxPlugin::getOutputDimensions(int index, nvinfer1::Dims YoloBoxPlugin::getOutputDimensions(
const nvinfer1::Dims* inputs, int index, const nvinfer1::Dims* inputs, int nb_input_dims) TRT_NOEXCEPT {
int nb_input_dims) {
const int anchor_num = anchors_.size() / 2; const int anchor_num = anchors_.size() / 2;
const int box_num = inputs[0].d[1] * inputs[0].d[2] * anchor_num; const int box_num = inputs[0].d[1] * inputs[0].d[2] * anchor_num;
...@@ -90,13 +91,15 @@ nvinfer1::Dims YoloBoxPlugin::getOutputDimensions(int index, ...@@ -90,13 +91,15 @@ nvinfer1::Dims YoloBoxPlugin::getOutputDimensions(int index,
return nvinfer1::Dims2(box_num, class_num_); return nvinfer1::Dims2(box_num, class_num_);
} }
bool YoloBoxPlugin::supportsFormat(nvinfer1::DataType type, bool YoloBoxPlugin::supportsFormat(
nvinfer1::TensorFormat format) const { nvinfer1::DataType type, nvinfer1::TensorFormat format) const TRT_NOEXCEPT {
return ((type == data_type_ || type == nvinfer1::DataType::kINT32) && return ((type == data_type_ || type == nvinfer1::DataType::kINT32) &&
format == nvinfer1::TensorFormat::kLINEAR); format == nvinfer1::TensorFormat::kLINEAR);
} }
size_t YoloBoxPlugin::getWorkspaceSize(int max_batch_size) const { return 0; } size_t YoloBoxPlugin::getWorkspaceSize(int max_batch_size) const TRT_NOEXCEPT {
return 0;
}
template <typename T> template <typename T>
__device__ inline T sigmoid(T x) { __device__ inline T sigmoid(T x) {
...@@ -219,7 +222,7 @@ __global__ void KeYoloBoxFw(const T* const input, const int* const imgsize, ...@@ -219,7 +222,7 @@ __global__ void KeYoloBoxFw(const T* const input, const int* const imgsize,
template <typename T> template <typename T>
int YoloBoxPlugin::enqueue_impl(int batch_size, const void* const* inputs, int YoloBoxPlugin::enqueue_impl(int batch_size, const void* const* inputs,
void** outputs, void* workspace, void* const* outputs, void* workspace,
cudaStream_t stream) { cudaStream_t stream) {
const int n = batch_size; const int n = batch_size;
const int h = input_h_; const int h = input_h_;
...@@ -247,7 +250,7 @@ int YoloBoxPlugin::enqueue(int batch_size, const void* const* inputs, ...@@ -247,7 +250,7 @@ int YoloBoxPlugin::enqueue(int batch_size, const void* const* inputs,
#else #else
void* const* outputs, void* workspace, void* const* outputs, void* workspace,
#endif #endif
cudaStream_t stream) { cudaStream_t stream) TRT_NOEXCEPT {
if (data_type_ == nvinfer1::DataType::kFLOAT) { if (data_type_ == nvinfer1::DataType::kFLOAT) {
return enqueue_impl<float>(batch_size, inputs, outputs, workspace, stream); return enqueue_impl<float>(batch_size, inputs, outputs, workspace, stream);
} else if (data_type_ == nvinfer1::DataType::kHALF) { } else if (data_type_ == nvinfer1::DataType::kHALF) {
...@@ -256,11 +259,11 @@ int YoloBoxPlugin::enqueue(int batch_size, const void* const* inputs, ...@@ -256,11 +259,11 @@ int YoloBoxPlugin::enqueue(int batch_size, const void* const* inputs,
assert("unsupported type."); assert("unsupported type.");
} }
int YoloBoxPlugin::initialize() { return 0; } int YoloBoxPlugin::initialize() TRT_NOEXCEPT { return 0; }
void YoloBoxPlugin::terminate() {} void YoloBoxPlugin::terminate() TRT_NOEXCEPT {}
size_t YoloBoxPlugin::getSerializationSize() const { size_t YoloBoxPlugin::getSerializationSize() const TRT_NOEXCEPT {
size_t serialize_size = 0; size_t serialize_size = 0;
serialize_size += SerializedSize(data_type_); serialize_size += SerializedSize(data_type_);
serialize_size += SerializedSize(anchors_); serialize_size += SerializedSize(anchors_);
...@@ -274,7 +277,7 @@ size_t YoloBoxPlugin::getSerializationSize() const { ...@@ -274,7 +277,7 @@ size_t YoloBoxPlugin::getSerializationSize() const {
return serialize_size; return serialize_size;
} }
void YoloBoxPlugin::serialize(void* buffer) const { void YoloBoxPlugin::serialize(void* buffer) const TRT_NOEXCEPT {
SerializeValue(&buffer, data_type_); SerializeValue(&buffer, data_type_);
SerializeValue(&buffer, anchors_); SerializeValue(&buffer, anchors_);
SerializeValue(&buffer, class_num_); SerializeValue(&buffer, class_num_);
...@@ -286,28 +289,30 @@ void YoloBoxPlugin::serialize(void* buffer) const { ...@@ -286,28 +289,30 @@ void YoloBoxPlugin::serialize(void* buffer) const {
SerializeValue(&buffer, input_w_); SerializeValue(&buffer, input_w_);
} }
void YoloBoxPlugin::destroy() {} void YoloBoxPlugin::destroy() TRT_NOEXCEPT {}
void YoloBoxPlugin::setPluginNamespace(const char* lib_namespace) { void YoloBoxPlugin::setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT {
namespace_ = std::string(lib_namespace); namespace_ = std::string(lib_namespace);
} }
const char* YoloBoxPlugin::getPluginNamespace() const { const char* YoloBoxPlugin::getPluginNamespace() const TRT_NOEXCEPT {
return namespace_.c_str(); return namespace_.c_str();
} }
nvinfer1::DataType YoloBoxPlugin::getOutputDataType( nvinfer1::DataType YoloBoxPlugin::getOutputDataType(
int index, const nvinfer1::DataType* input_type, int nb_inputs) const { int index, const nvinfer1::DataType* input_type,
int nb_inputs) const TRT_NOEXCEPT {
return input_type[0]; return input_type[0];
} }
bool YoloBoxPlugin::isOutputBroadcastAcrossBatch(int output_index, bool YoloBoxPlugin::isOutputBroadcastAcrossBatch(
const bool* input_is_broadcast, int output_index, const bool* input_is_broadcast,
int nb_inputs) const { int nb_inputs) const TRT_NOEXCEPT {
return false; return false;
} }
bool YoloBoxPlugin::canBroadcastInputAcrossBatch(int input_index) const { bool YoloBoxPlugin::canBroadcastInputAcrossBatch(int input_index) const
TRT_NOEXCEPT {
return false; return false;
} }
...@@ -317,9 +322,9 @@ void YoloBoxPlugin::configurePlugin( ...@@ -317,9 +322,9 @@ void YoloBoxPlugin::configurePlugin(
const nvinfer1::DataType* input_types, const nvinfer1::DataType* input_types,
const nvinfer1::DataType* output_types, const bool* input_is_broadcast, const nvinfer1::DataType* output_types, const bool* input_is_broadcast,
const bool* output_is_broadcast, nvinfer1::PluginFormat float_format, const bool* output_is_broadcast, nvinfer1::PluginFormat float_format,
int max_batct_size) {} int max_batct_size) TRT_NOEXCEPT {}
nvinfer1::IPluginV2Ext* YoloBoxPlugin::clone() const { nvinfer1::IPluginV2Ext* YoloBoxPlugin::clone() const TRT_NOEXCEPT {
return new YoloBoxPlugin(data_type_, anchors_, class_num_, conf_thresh_, return new YoloBoxPlugin(data_type_, anchors_, class_num_, conf_thresh_,
downsample_ratio_, clip_bbox_, scale_x_y_, input_h_, downsample_ratio_, clip_bbox_, scale_x_y_, input_h_,
input_w_); input_w_);
...@@ -327,26 +332,30 @@ nvinfer1::IPluginV2Ext* YoloBoxPlugin::clone() const { ...@@ -327,26 +332,30 @@ nvinfer1::IPluginV2Ext* YoloBoxPlugin::clone() const {
YoloBoxPluginCreator::YoloBoxPluginCreator() {} YoloBoxPluginCreator::YoloBoxPluginCreator() {}
void YoloBoxPluginCreator::setPluginNamespace(const char* lib_namespace) { void YoloBoxPluginCreator::setPluginNamespace(const char* lib_namespace)
TRT_NOEXCEPT {
namespace_ = std::string(lib_namespace); namespace_ = std::string(lib_namespace);
} }
const char* YoloBoxPluginCreator::getPluginNamespace() const { const char* YoloBoxPluginCreator::getPluginNamespace() const TRT_NOEXCEPT {
return namespace_.c_str(); return namespace_.c_str();
} }
const char* YoloBoxPluginCreator::getPluginName() const { const char* YoloBoxPluginCreator::getPluginName() const TRT_NOEXCEPT {
return "yolo_box_plugin"; return "yolo_box_plugin";
} }
const char* YoloBoxPluginCreator::getPluginVersion() const { return "1"; } const char* YoloBoxPluginCreator::getPluginVersion() const TRT_NOEXCEPT {
return "1";
}
const nvinfer1::PluginFieldCollection* YoloBoxPluginCreator::getFieldNames() { const nvinfer1::PluginFieldCollection* YoloBoxPluginCreator::getFieldNames()
TRT_NOEXCEPT {
return &field_collection_; return &field_collection_;
} }
nvinfer1::IPluginV2Ext* YoloBoxPluginCreator::createPlugin( nvinfer1::IPluginV2Ext* YoloBoxPluginCreator::createPlugin(
const char* name, const nvinfer1::PluginFieldCollection* fc) { const char* name, const nvinfer1::PluginFieldCollection* fc) TRT_NOEXCEPT {
const nvinfer1::PluginField* fields = fc->fields; const nvinfer1::PluginField* fields = fc->fields;
int type_id = -1; int type_id = -1;
...@@ -392,7 +401,8 @@ nvinfer1::IPluginV2Ext* YoloBoxPluginCreator::createPlugin( ...@@ -392,7 +401,8 @@ nvinfer1::IPluginV2Ext* YoloBoxPluginCreator::createPlugin(
} }
nvinfer1::IPluginV2Ext* YoloBoxPluginCreator::deserializePlugin( nvinfer1::IPluginV2Ext* YoloBoxPluginCreator::deserializePlugin(
const char* name, const void* serial_data, size_t serial_length) { const char* name, const void* serial_data,
size_t serial_length) TRT_NOEXCEPT {
auto plugin = new YoloBoxPlugin(serial_data, serial_length); auto plugin = new YoloBoxPlugin(serial_data, serial_length);
plugin->setPluginNamespace(namespace_.c_str()); plugin->setPluginNamespace(namespace_.c_str());
return plugin; return plugin;
......
paddle/fluid/inference/tensorrt/plugin/yolo_box_op_plugin.h
浏览文件 @ b5aab4f0
...@@ -35,38 +35,39 @@ class YoloBoxPlugin : public nvinfer1::IPluginV2Ext { ...@@ -35,38 +35,39 @@ class YoloBoxPlugin : public nvinfer1::IPluginV2Ext {
YoloBoxPlugin(const void* data, size_t length); YoloBoxPlugin(const void* data, size_t length);
~YoloBoxPlugin() override; ~YoloBoxPlugin() override;
const char* getPluginType() const override; const char* getPluginType() const TRT_NOEXCEPT override;
const char* getPluginVersion() const override; const char* getPluginVersion() const TRT_NOEXCEPT override;
int getNbOutputs() const override; int getNbOutputs() const TRT_NOEXCEPT override;
nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs,
int nb_input_dims) override; int nb_input_dims) TRT_NOEXCEPT override;
bool supportsFormat(nvinfer1::DataType type, bool supportsFormat(nvinfer1::DataType type, nvinfer1::TensorFormat format)
nvinfer1::TensorFormat format) const override; const TRT_NOEXCEPT override;
size_t getWorkspaceSize(int max_batch_size) const override; size_t getWorkspaceSize(int max_batch_size) const TRT_NOEXCEPT override;
#if IS_TRT_VERSION_LT(8000) #if IS_TRT_VERSION_LT(8000)
int enqueue(int batch_size, const void* const* inputs, void** outputs, int enqueue(int batch_size, const void* const* inputs, void** outputs,
#else #else
int enqueue(int batch_size, const void* const* inputs, void* const* outputs, int enqueue(int batch_size, const void* const* inputs, void* const* outputs,
#endif #endif
void* workspace, cudaStream_t stream) override; void* workspace, cudaStream_t stream) TRT_NOEXCEPT override;
template <typename T> template <typename T>
int enqueue_impl(int batch_size, const void* const* inputs, void** outputs, int enqueue_impl(int batch_size, const void* const* inputs,
void* workspace, cudaStream_t stream); void* const* outputs, void* workspace, cudaStream_t stream);
int initialize() override; int initialize() TRT_NOEXCEPT override;
void terminate() override; void terminate() TRT_NOEXCEPT override;
size_t getSerializationSize() const override; size_t getSerializationSize() const TRT_NOEXCEPT override;
void serialize(void* buffer) const override; void serialize(void* buffer) const TRT_NOEXCEPT override;
void destroy() override; void destroy() TRT_NOEXCEPT override;
void setPluginNamespace(const char* lib_namespace) override; void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override;
const char* getPluginNamespace() const override; const char* getPluginNamespace() const TRT_NOEXCEPT override;
nvinfer1::DataType getOutputDataType(int index, nvinfer1::DataType getOutputDataType(
const nvinfer1::DataType* input_type, int index, const nvinfer1::DataType* input_type,
int nb_inputs) const override; int nb_inputs) const TRT_NOEXCEPT override;
bool isOutputBroadcastAcrossBatch(int output_index, bool isOutputBroadcastAcrossBatch(int output_index,
const bool* input_is_broadcast, const bool* input_is_broadcast,
int nb_inputs) const override; int nb_inputs) const TRT_NOEXCEPT override;
bool canBroadcastInputAcrossBatch(int input_index) const override; bool canBroadcastInputAcrossBatch(int input_index) const
TRT_NOEXCEPT override;
void configurePlugin(const nvinfer1::Dims* input_dims, int nb_inputs, void configurePlugin(const nvinfer1::Dims* input_dims, int nb_inputs,
const nvinfer1::Dims* output_dims, int nb_outputs, const nvinfer1::Dims* output_dims, int nb_outputs,
const nvinfer1::DataType* input_types, const nvinfer1::DataType* input_types,
...@@ -74,8 +75,8 @@ class YoloBoxPlugin : public nvinfer1::IPluginV2Ext { ...@@ -74,8 +75,8 @@ class YoloBoxPlugin : public nvinfer1::IPluginV2Ext {
const bool* input_is_broadcast, const bool* input_is_broadcast,
const bool* output_is_broadcast, const bool* output_is_broadcast,
nvinfer1::PluginFormat float_format, nvinfer1::PluginFormat float_format,
int max_batct_size) override; int max_batct_size) TRT_NOEXCEPT override;
nvinfer1::IPluginV2Ext* clone() const override; nvinfer1::IPluginV2Ext* clone() const TRT_NOEXCEPT override;
private: private:
nvinfer1::DataType data_type_; nvinfer1::DataType data_type_;
...@@ -96,17 +97,18 @@ class YoloBoxPluginCreator : public nvinfer1::IPluginCreator { ...@@ -96,17 +97,18 @@ class YoloBoxPluginCreator : public nvinfer1::IPluginCreator {
YoloBoxPluginCreator(); YoloBoxPluginCreator();
~YoloBoxPluginCreator() override = default; ~YoloBoxPluginCreator() override = default;
void setPluginNamespace(const char* lib_namespace) override; void setPluginNamespace(const char* lib_namespace) TRT_NOEXCEPT override;
const char* getPluginNamespace() const override; const char* getPluginNamespace() const TRT_NOEXCEPT override;
const char* getPluginName() const override; const char* getPluginName() const TRT_NOEXCEPT override;
const char* getPluginVersion() const override; const char* getPluginVersion() const TRT_NOEXCEPT override;
const nvinfer1::PluginFieldCollection* getFieldNames() override; const nvinfer1::PluginFieldCollection* getFieldNames() TRT_NOEXCEPT override;
nvinfer1::IPluginV2Ext* createPlugin( nvinfer1::IPluginV2Ext* createPlugin(
const char* name, const nvinfer1::PluginFieldCollection* fc) override; const char* name,
nvinfer1::IPluginV2Ext* deserializePlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) TRT_NOEXCEPT override;
const void* serial_data, nvinfer1::IPluginV2Ext* deserializePlugin(
size_t serial_length) override; const char* name, const void* serial_data,
size_t serial_length) TRT_NOEXCEPT override;
private: private:
std::string namespace_; std::string namespace_;
......
paddle/fluid/inference/tensorrt/trt_int8_calibrator.cc
浏览文件 @ b5aab4f0
...@@ -22,7 +22,7 @@ namespace inference { ...@@ -22,7 +22,7 @@ namespace inference {
namespace tensorrt { namespace tensorrt {
// set the batch size before constructing the thread to execute engine // set the batch size before constructing the thread to execute engine
int TRTInt8Calibrator::getBatchSize() const { return batch_size_; } int TRTInt8Calibrator::getBatchSize() const TRT_NOEXCEPT { return batch_size_; }
TRTInt8Calibrator::TRTInt8Calibrator( TRTInt8Calibrator::TRTInt8Calibrator(
const std::unordered_map<std::string, size_t>& buffers, int batch_size, const std::unordered_map<std::string, size_t>& buffers, int batch_size,
...@@ -95,7 +95,7 @@ bool TRTInt8Calibrator::setBatch( ...@@ -95,7 +95,7 @@ bool TRTInt8Calibrator::setBatch(
} }
bool TRTInt8Calibrator::getBatch(void** bindings, const char** names, bool TRTInt8Calibrator::getBatch(void** bindings, const char** names,
int num_bindings) { int num_bindings) TRT_NOEXCEPT {
VLOG(4) << "get batch: " << engine_name_; VLOG(4) << "get batch: " << engine_name_;
std::unique_lock<std::mutex> lk(mut_); std::unique_lock<std::mutex> lk(mut_);
// The consumer has just finished processing a data. // The consumer has just finished processing a data.
...@@ -131,14 +131,15 @@ void TRTInt8Calibrator::setDone() { ...@@ -131,14 +131,15 @@ void TRTInt8Calibrator::setDone() {
cond_.notify_all(); cond_.notify_all();
} }
const void* TRTInt8Calibrator::readCalibrationCache(size_t& length) { const void* TRTInt8Calibrator::readCalibrationCache(size_t& length)
TRT_NOEXCEPT {
if (calibration_table_.empty()) return nullptr; if (calibration_table_.empty()) return nullptr;
length = calibration_table_.size(); length = calibration_table_.size();
return calibration_table_.data(); return calibration_table_.data();
} }
void TRTInt8Calibrator::writeCalibrationCache(const void* ptr, void TRTInt8Calibrator::writeCalibrationCache(const void* ptr,
std::size_t length) { std::size_t length) TRT_NOEXCEPT {
calibration_table_ = std::string((const char*)ptr, length); calibration_table_ = std::string((const char*)ptr, length);
VLOG(4) << "Got calibration data for " << engine_name_ << " " << ptr VLOG(4) << "Got calibration data for " << engine_name_ << " " << ptr
<< " length=" << length; << " length=" << length;
......
paddle/fluid/inference/tensorrt/trt_int8_calibrator.h
浏览文件 @ b5aab4f0
...@@ -43,17 +43,18 @@ struct TRTInt8Calibrator : public nvinfer1::IInt8EntropyCalibrator2 { ...@@ -43,17 +43,18 @@ struct TRTInt8Calibrator : public nvinfer1::IInt8EntropyCalibrator2 {
explicit TRTInt8Calibrator(const std::string& calibration_data); explicit TRTInt8Calibrator(const std::string& calibration_data);
~TRTInt8Calibrator(); ~TRTInt8Calibrator();
int getBatchSize() const override; int getBatchSize() const TRT_NOEXCEPT override;
bool getBatch(void* bindings[], const char* names[], bool getBatch(void* bindings[], const char* names[],
int num_bindings) override; int num_bindings) TRT_NOEXCEPT override;
bool setBatch(const std::unordered_map<std::string, void*>& data); bool setBatch(const std::unordered_map<std::string, void*>& data);
void setDone(); void setDone();
void waitAndSetDone(); void waitAndSetDone();
const void* readCalibrationCache(std::size_t& length) override; const void* readCalibrationCache(std::size_t& length) TRT_NOEXCEPT override;
void writeCalibrationCache(const void* ptr, std::size_t length) override; void writeCalibrationCache(const void* ptr,
std::size_t length) TRT_NOEXCEPT override;
const std::string& getCalibrationTableAsString() { const std::string& getCalibrationTableAsString() {
return calibration_table_; return calibration_table_;
} }
......
Markdown is supported
0% 或 .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑！
想要评论请 注册 或