diff --git a/include/tkDNN/pluginsRT/ActivationLeakyRT.h b/include/tkDNN/pluginsRT/ActivationLeakyRT.h index ae395ef..1d98a59 100644 --- a/include/tkDNN/pluginsRT/ActivationLeakyRT.h +++ b/include/tkDNN/pluginsRT/ActivationLeakyRT.h @@ -27,8 +27,12 @@ namespace nvinfer1 { size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override; +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, void const *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override; diff --git a/include/tkDNN/pluginsRT/ActivationLogisticRT.h b/include/tkDNN/pluginsRT/ActivationLogisticRT.h index 9646efc..d972752 100644 --- a/include/tkDNN/pluginsRT/ActivationLogisticRT.h +++ b/include/tkDNN/pluginsRT/ActivationLogisticRT.h @@ -28,8 +28,13 @@ namespace nvinfer1 { size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override; +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif + size_t getSerializationSize() const NOEXCEPT override ; diff --git a/include/tkDNN/pluginsRT/ActivationMishRT.h b/include/tkDNN/pluginsRT/ActivationMishRT.h index 3c9f15a..5b966cd 100644 --- a/include/tkDNN/pluginsRT/ActivationMishRT.h +++ b/include/tkDNN/pluginsRT/ActivationMishRT.h @@ -26,8 +26,11 @@ namespace nvinfer1 { void terminate() NOEXCEPT override ; size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; - +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT override ; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override ; diff --git a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h index b5c186c..1830945 100644 --- a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h +++ b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h @@ -25,8 +25,11 @@ namespace nvinfer1 { void terminate() NOEXCEPT override ; size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; - +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT override ; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override ; diff --git a/include/tkDNN/pluginsRT/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h index 1cbf55d..e0541ef 100644 --- a/include/tkDNN/pluginsRT/DeformableConvRT.h +++ b/include/tkDNN/pluginsRT/DeformableConvRT.h @@ -33,9 +33,12 @@ namespace nvinfer1 { void terminate() NOEXCEPT override ; size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; - +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override ; diff --git a/include/tkDNN/pluginsRT/FlattenConcatRT.h b/include/tkDNN/pluginsRT/FlattenConcatRT.h index 7d655d3..b1e5327 100644 --- a/include/tkDNN/pluginsRT/FlattenConcatRT.h +++ b/include/tkDNN/pluginsRT/FlattenConcatRT.h @@ -25,7 +25,11 @@ namespace nvinfer1 { size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override ; diff --git a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h index 14f714e..723bed6 100644 --- a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h +++ b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h @@ -28,8 +28,12 @@ namespace nvinfer1 { size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override ; diff --git a/include/tkDNN/pluginsRT/RegionRT.h b/include/tkDNN/pluginsRT/RegionRT.h index cdcb515..400f2f0 100644 --- a/include/tkDNN/pluginsRT/RegionRT.h +++ b/include/tkDNN/pluginsRT/RegionRT.h @@ -30,8 +30,12 @@ namespace nvinfer1 { size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override ; diff --git a/include/tkDNN/pluginsRT/ReorgRT.h b/include/tkDNN/pluginsRT/ReorgRT.h index a780615..ced6b9c 100644 --- a/include/tkDNN/pluginsRT/ReorgRT.h +++ b/include/tkDNN/pluginsRT/ReorgRT.h @@ -27,8 +27,12 @@ namespace nvinfer1 { size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override; +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override; diff --git a/include/tkDNN/pluginsRT/ReshapeRT.h b/include/tkDNN/pluginsRT/ReshapeRT.h index 8d3fdb5..a0cc718 100644 --- a/include/tkDNN/pluginsRT/ReshapeRT.h +++ b/include/tkDNN/pluginsRT/ReshapeRT.h @@ -31,7 +31,11 @@ namespace nvinfer1 { size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override ; diff --git a/include/tkDNN/pluginsRT/ResizeLayerRT.h b/include/tkDNN/pluginsRT/ResizeLayerRT.h index 7316048..a446698 100644 --- a/include/tkDNN/pluginsRT/ResizeLayerRT.h +++ b/include/tkDNN/pluginsRT/ResizeLayerRT.h @@ -29,8 +29,13 @@ namespace nvinfer1 { size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif + size_t getSerializationSize() const NOEXCEPT override ; diff --git a/include/tkDNN/pluginsRT/RouteRT.h b/include/tkDNN/pluginsRT/RouteRT.h index 2b9bc67..499b9da 100644 --- a/include/tkDNN/pluginsRT/RouteRT.h +++ b/include/tkDNN/pluginsRT/RouteRT.h @@ -29,7 +29,11 @@ namespace nvinfer1 { size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT override ; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override ; diff --git a/include/tkDNN/pluginsRT/ShortcutRT.h b/include/tkDNN/pluginsRT/ShortcutRT.h index 7069180..29afc85 100644 --- a/include/tkDNN/pluginsRT/ShortcutRT.h +++ b/include/tkDNN/pluginsRT/ShortcutRT.h @@ -32,8 +32,12 @@ namespace nvinfer1 { size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override; +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override; diff --git a/include/tkDNN/pluginsRT/UpsampleRT.h b/include/tkDNN/pluginsRT/UpsampleRT.h index 030fe5a..09f3547 100644 --- a/include/tkDNN/pluginsRT/UpsampleRT.h +++ b/include/tkDNN/pluginsRT/UpsampleRT.h @@ -30,8 +30,12 @@ namespace nvinfer1 { size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override; +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override; diff --git a/include/tkDNN/pluginsRT/YoloRT.h b/include/tkDNN/pluginsRT/YoloRT.h index 0d895e0..9479f9f 100644 --- a/include/tkDNN/pluginsRT/YoloRT.h +++ b/include/tkDNN/pluginsRT/YoloRT.h @@ -34,8 +34,13 @@ namespace nvinfer1 { size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override; + +#if NV_TENSORRT_MAJOR > 7 int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override; +#elif NV_TENSORRT_MAJOR == 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override; +#endif size_t getSerializationSize() const NOEXCEPT override; diff --git a/src/pluginsRT/ActivationLeakyRT.cpp b/src/pluginsRT/ActivationLeakyRT.cpp index 77063f6..2a9bf33 100644 --- a/src/pluginsRT/ActivationLeakyRT.cpp +++ b/src/pluginsRT/ActivationLeakyRT.cpp @@ -42,6 +42,7 @@ size_t ActivationLeakyRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } +#if NV_TENSORRT_MAJOR > 7 int ActivationLeakyRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { activationLEAKYForward( @@ -51,6 +52,17 @@ int ActivationLeakyRT::enqueue(int batchSize, const void *const *inputs, void *c return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t ActivationLeakyRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, + cudaStream_t stream) { + activationLEAKYForward( + (dnnType *) reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize * size, slope, + stream); + return 0; +} +#endif + size_t ActivationLeakyRT::getSerializationSize() const NOEXCEPT { return 1 * sizeof(int) + 1 * sizeof(float); @@ -94,8 +106,7 @@ IPluginV2* ActivationLeakyRT::clone() const NOEXCEPT { } ActivationLeakyRTPluginCreator::ActivationLeakyRTPluginCreator() { - mPluginAttributes.emplace_back( - PluginField("slope", nullptr, PluginFieldType::kFLOAT32, 1)); + mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); mFC.fields = mPluginAttributes.data(); } diff --git a/src/pluginsRT/ActivationLogisticRT.cpp b/src/pluginsRT/ActivationLogisticRT.cpp index 9966708..a3c4ae7 100644 --- a/src/pluginsRT/ActivationLogisticRT.cpp +++ b/src/pluginsRT/ActivationLogisticRT.cpp @@ -38,13 +38,21 @@ void ActivationLogisticRT::terminate() NOEXCEPT {} size_t ActivationLogisticRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } - +#if NV_TENSORRT_MAJOR > 7 int ActivationLogisticRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { activationLOGISTICForward((dnnType *) reinterpret_cast(inputs[0]), reinterpret_cast(outputs[0]), batchSize * size, stream); return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t ActivationLogisticRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, + cudaStream_t stream) { + activationLOGISTICForward((dnnType *) reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize * size, stream); + return 0; +} +#endif size_t ActivationLogisticRT::getSerializationSize() const NOEXCEPT { return 1 * sizeof(int); diff --git a/src/pluginsRT/ActivationMishRT.cpp b/src/pluginsRT/ActivationMishRT.cpp index 7797df9..d42b90a 100644 --- a/src/pluginsRT/ActivationMishRT.cpp +++ b/src/pluginsRT/ActivationMishRT.cpp @@ -38,12 +38,22 @@ void ActivationMishRT::terminate() NOEXCEPT {} size_t ActivationMishRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } +#if NV_TENSORRT_MAJOR > 7 int ActivationMishRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { activationMishForward((dnnType *) reinterpret_cast(inputs[0]), reinterpret_cast(outputs[0]), batchSize * size, stream); return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t ActivationMishRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, + cudaStream_t stream) { + activationMishForward((dnnType *) reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize * size, stream); + return 0; +} +#endif + size_t ActivationMishRT::getSerializationSize() const NOEXCEPT { return 1 * sizeof(int); } @@ -80,6 +90,8 @@ IPluginV2 *ActivationMishRT::clone() const NOEXCEPT { return p; } + + ActivationMishRTPluginCreator::ActivationMishRTPluginCreator() { mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); diff --git a/src/pluginsRT/ActivationReLUCeilingRT.cpp b/src/pluginsRT/ActivationReLUCeilingRT.cpp index 80b5a74..c97d4ec 100644 --- a/src/pluginsRT/ActivationReLUCeilingRT.cpp +++ b/src/pluginsRT/ActivationReLUCeilingRT.cpp @@ -40,11 +40,20 @@ size_t ActivationReLUCeiling::getWorkspaceSize(int maxBatchSize) const NOEXCEPT return 0; } +#if NV_TENSORRT_MAJOR > 7 int ActivationReLUCeiling::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT { activationReLUCeilingForward((dnnType *) reinterpret_cast(inputs[0]), reinterpret_cast(outputs[0]), batchSize * size, ceiling, stream); return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t ActivationReLUCeiling::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, + cudaStream_t stream) { + activationReLUCeilingForward((dnnType *) reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize * size, ceiling, stream); + return 0; +} +#endif size_t ActivationReLUCeiling::getSerializationSize() const NOEXCEPT { return 1 * sizeof(int) + 1 * sizeof(float); @@ -86,7 +95,7 @@ void ActivationReLUCeiling::setPluginNamespace(const char *pluginNamespace) NOEX } ActivationReLUCeilingPluginCreator::ActivationReLUCeilingPluginCreator() { - mPluginAttributes.emplace_back(PluginField("ceiling", nullptr, PluginFieldType::kFLOAT32, 1)); + mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); mFC.fields = mPluginAttributes.data(); } diff --git a/src/pluginsRT/DeformableConvRT.cpp b/src/pluginsRT/DeformableConvRT.cpp index 417af8e..62f7555 100644 --- a/src/pluginsRT/DeformableConvRT.cpp +++ b/src/pluginsRT/DeformableConvRT.cpp @@ -137,6 +137,7 @@ void DeformableConvRT::terminate() NOEXCEPT {} size_t DeformableConvRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {return 0;} +#if NV_TENSORRT_MAJOR > 7 int DeformableConvRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); @@ -165,6 +166,36 @@ int DeformableConvRT::enqueue(int batchSize, const void *const *inputs, void *co } return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t DeformableConvRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, + cudaStream_t stream) { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *output_conv = (dnnType*)reinterpret_cast(inputs[1]); + + // split conv2d outputs into offset to mask + for(int b=0; b(outputs[0]), ones_d2, + kh, kw, + sh, sw, + ph, pw, + 1, 1, + deformableGroup, b, + i_n, i_c, i_h, i_w, + o_n, o_c, o_h, o_w, + chunk_dim); + } + return 0; +} +#endif size_t DeformableConvRT::getSerializationSize() const NOEXCEPT { return 16 * sizeof(int) + chunk_dim * 3 * sizeof(dnnType) + (i_c * o_c * kh * kw * 1 ) * sizeof(dnnType) + @@ -251,25 +282,8 @@ IPluginV2 *DeformableConvRT::clone() const NOEXCEPT { return p; } - DeformableConvRTPluginCreator::DeformableConvRTPluginCreator() { - mPluginAttributes.emplace_back(PluginField("chunk_dim",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("kh",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("kw",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("sh",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("sw",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("ph",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("pw",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("deformableGroup",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("i_n",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("i_c",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("i_h",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("i_w",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_n",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_c",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_h",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_w",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("defRT",nullptr,PluginFieldType::kUNKNOWN,1)); + mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); mFC.fields = mPluginAttributes.data(); } diff --git a/src/pluginsRT/FlattenConcatRT.cpp b/src/pluginsRT/FlattenConcatRT.cpp index 51cebc5..6687b58 100644 --- a/src/pluginsRT/FlattenConcatRT.cpp +++ b/src/pluginsRT/FlattenConcatRT.cpp @@ -54,6 +54,7 @@ size_t FlattenConcatRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } +#if NV_TENSORRT_MAJOR > 7 int FlattenConcatRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); @@ -69,6 +70,24 @@ int FlattenConcatRT::enqueue(int batchSize, const void *const *inputs, void *con } return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t FlattenConcatRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, + cudaStream_t stream) { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); + + checkERROR( cublasSetStream(handle, stream) ); + for(int i=0; i 7 int MaxPoolFixedSizeRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); @@ -61,6 +62,16 @@ int MaxPoolFixedSizeRT::enqueue(int batchSize, const void *const *inputs, void * MaxPoolingForward(srcData, dstData, batchSize, this->c, this->h, this->w, this->stride_H, this->stride_W, this->winSize, this->padding, stream); return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t MaxPoolFixedSizeRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, + cudaStream_t stream) { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + MaxPoolingForward(srcData, dstData, batchSize, this->c, this->h, this->w, this->stride_H, this->stride_W, this->winSize, this->padding, stream); + return 0; +} +#endif + size_t MaxPoolFixedSizeRT::getSerializationSize() const NOEXCEPT { return 8*sizeof(int); @@ -110,16 +121,8 @@ IPluginV2 *MaxPoolFixedSizeRT::clone() const NOEXCEPT { return p; } - MaxPoolFixedSizeRTPluginCreator::MaxPoolFixedSizeRTPluginCreator() { - mPluginAttributes.emplace_back(PluginField("c",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("h",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("w",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("n",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("stride_H",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("stride_W",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("winSize",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("padding",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); mFC.fields = mPluginAttributes.data(); } diff --git a/src/pluginsRT/RegionRT.cpp b/src/pluginsRT/RegionRT.cpp index de1dd3b..552db3e 100644 --- a/src/pluginsRT/RegionRT.cpp +++ b/src/pluginsRT/RegionRT.cpp @@ -43,6 +43,7 @@ void RegionRT::terminate() NOEXCEPT {} size_t RegionRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } +#if NV_TENSORRT_MAJOR > 7 int RegionRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); @@ -68,6 +69,32 @@ int RegionRT::enqueue(int batchSize, const void *const *inputs, void *const *out return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t RegionRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + + checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); + + for (int b = 0; b < batchSize; ++b){ + for(int n = 0; n < num; ++n){ + int index = entry_index(b, n*w*h, 0); + activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); + + index = entry_index(b, n*w*h, coords); + activationLOGISTICForward(srcData + index, dstData + index, w*h, stream); + } + } + + //softmax start + int index = entry_index(0, 0, coords + 1); + softmaxForward( srcData + index, classes, batchSize*num, + (c*h*w)/num, + w*h, 1, w*h, 1, dstData + index, stream); + + return 0; +} +#endif size_t RegionRT::getSerializationSize() const NOEXCEPT { return 6*sizeof(int); @@ -112,11 +139,8 @@ IPluginV2 *RegionRT::clone() const NOEXCEPT { return p; } - RegionRTPluginCreator::RegionRTPluginCreator() { - mPluginAttributes.emplace_back(PluginField("classes",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("coords",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("num",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); mFC.fields = mPluginAttributes.data(); } diff --git a/src/pluginsRT/ReorgRT.cpp b/src/pluginsRT/ReorgRT.cpp index d3dd12a..c2b0083 100644 --- a/src/pluginsRT/ReorgRT.cpp +++ b/src/pluginsRT/ReorgRT.cpp @@ -43,12 +43,22 @@ size_t ReorgRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } +#if NV_TENSORRT_MAJOR > 7 int ReorgRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT { reorgForward((dnnType*)reinterpret_cast(inputs[0]), reinterpret_cast(outputs[0]), batchSize, c, h, w, stride, stream); return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t ReorgRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) { + reorgForward((dnnType*)reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), + batchSize, c, h, w, stride, stream); + return 0; +} +#endif + size_t ReorgRT::getSerializationSize() const NOEXCEPT { return 4*sizeof(int); @@ -93,9 +103,8 @@ IPluginV2 *ReorgRT::clone() const NOEXCEPT { return p; } - ReorgRTPluginCreator::ReorgRTPluginCreator() { - mPluginAttributes.emplace_back(PluginField("stride",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); mFC.fields = mPluginAttributes.data(); } diff --git a/src/pluginsRT/ReshapeRT.cpp b/src/pluginsRT/ReshapeRT.cpp index 5e23ab3..127e244 100644 --- a/src/pluginsRT/ReshapeRT.cpp +++ b/src/pluginsRT/ReshapeRT.cpp @@ -43,6 +43,7 @@ size_t ReshapeRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } +#if NV_TENSORRT_MAJOR > 7 int ReshapeRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); @@ -51,6 +52,16 @@ int ReshapeRT::enqueue(int batchSize, const void *const *inputs, void *const *ou checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t ReshapeRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + + checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); + return 0; +} +#endif + size_t ReshapeRT::getSerializationSize() const NOEXCEPT { return 4*sizeof(int); @@ -96,9 +107,8 @@ IPluginV2 *ReshapeRT::clone() const NOEXCEPT { return p; } - ReshapeRTPluginCreator::ReshapeRTPluginCreator() { - mPluginAttributes.emplace_back(PluginField("new_dim",nullptr,PluginFieldType::kUNKNOWN,1)); + mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); mFC.fields = mPluginAttributes.data(); } diff --git a/src/pluginsRT/ResizeLayerRT.cpp b/src/pluginsRT/ResizeLayerRT.cpp index 12a3ff9..dd8c912 100644 --- a/src/pluginsRT/ResizeLayerRT.cpp +++ b/src/pluginsRT/ResizeLayerRT.cpp @@ -47,6 +47,7 @@ void ResizeLayerRT::terminate() NOEXCEPT {} size_t ResizeLayerRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } +#if NV_TENSORRT_MAJOR > 7 int ResizeLayerRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { resizeForward((dnnType*)reinterpret_cast(inputs[0]), @@ -54,6 +55,15 @@ int ResizeLayerRT::enqueue(int batchSize, const void *const *inputs, void *const batchSize, i_c, i_h, i_w, o_c, o_h, o_w, stream); return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t ResizeLayerRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, + cudaStream_t stream) { + resizeForward((dnnType*)reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), + batchSize, i_c, i_h, i_w, o_c, o_h, o_w, stream); + return 0; +} +#endif size_t ResizeLayerRT::getSerializationSize() const NOEXCEPT { return 6*sizeof(int); @@ -101,11 +111,8 @@ IPluginV2 *ResizeLayerRT::clone() const NOEXCEPT { return p; } - ResizeLayerRTPluginCreator::ResizeLayerRTPluginCreator() { - mPluginAttributes.emplace_back(PluginField("o_c",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_h",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_w",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); mFC.fields = mPluginAttributes.data(); } diff --git a/src/pluginsRT/RouteRT.cpp b/src/pluginsRT/RouteRT.cpp index 87233a5..3218622 100644 --- a/src/pluginsRT/RouteRT.cpp +++ b/src/pluginsRT/RouteRT.cpp @@ -60,6 +60,7 @@ size_t RouteRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } +#if NV_TENSORRT_MAJOR > 7 int RouteRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { dnnType *dstData = reinterpret_cast(outputs[0]); @@ -75,6 +76,22 @@ int RouteRT::enqueue(int batchSize, const void *const *inputs, void *const *outp } return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t RouteRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) { + dnnType *dstData = reinterpret_cast(outputs[0]); + for(int b=0; b(inputs[i]); + int in_dim = c_in[i]*h*w; + int part_in_dim = in_dim / this->groups; + checkCuda( cudaMemcpyAsync(dstData + b*c*w*h + offset, input + b*c*w*h*groups + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) ); + offset += part_in_dim; + } + } + return 0; +} +#endif size_t RouteRT::getSerializationSize() const NOEXCEPT { return (6+MAX_INPUTS)*sizeof(int); @@ -124,10 +141,8 @@ IPluginV2 *RouteRT::clone() const NOEXCEPT { return p; } - RouteRTPluginCreator::RouteRTPluginCreator() { - mPluginAttributes.emplace_back(PluginField("groups",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("group_id",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); mFC.fields = mPluginAttributes.data(); } diff --git a/src/pluginsRT/ShortcutRT.cpp b/src/pluginsRT/ShortcutRT.cpp index eb35bd6..3cedffa 100644 --- a/src/pluginsRT/ShortcutRT.cpp +++ b/src/pluginsRT/ShortcutRT.cpp @@ -50,6 +50,7 @@ void ShortcutRT::terminate() NOEXCEPT {} size_t ShortcutRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } +#if NV_TENSORRT_MAJOR > 7 int ShortcutRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); @@ -61,6 +62,20 @@ int ShortcutRT::enqueue(int batchSize, const void *const *inputs, void *const *o return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t ShortcutRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, + cudaStream_t stream) { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *srcDataBack = (dnnType*)reinterpret_cast(inputs[1]); + dnnType *dstData = reinterpret_cast(outputs[0]); + + checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); + shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, bc, bh, bw, 1, mul, stream); + + return 0; +} +#endif + size_t ShortcutRT::getSerializationSize() const NOEXCEPT { return 6*sizeof(int) + sizeof(bool); @@ -108,10 +123,8 @@ IPluginV2 *ShortcutRT::clone() const NOEXCEPT { return p; } - ShortcutRTPluginCreator::ShortcutRTPluginCreator() { - mPluginAttributes.emplace_back(PluginField("bDim",nullptr,PluginFieldType::kUNKNOWN,1)); - mPluginAttributes.emplace_back(PluginField("mul",nullptr,PluginFieldType::kUNKNOWN,1)); + mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); mFC.fields = mPluginAttributes.data(); } diff --git a/src/pluginsRT/UpsampleRT.cpp b/src/pluginsRT/UpsampleRT.cpp index 044405b..d08a3ca 100644 --- a/src/pluginsRT/UpsampleRT.cpp +++ b/src/pluginsRT/UpsampleRT.cpp @@ -44,6 +44,7 @@ size_t UpsampleRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } +#if NV_TENSORRT_MAJOR > 7 int UpsampleRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); @@ -53,6 +54,17 @@ int UpsampleRT::enqueue(int batchSize, const void *const *inputs, void *const *o upsampleForward(srcData, dstData, batchSize, c, h, w, stride, 1, 1, stream); return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t UpsampleRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, + cudaStream_t stream) { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + + fill(dstData, batchSize*c*h*w*stride*stride, 0.0, stream); + upsampleForward(srcData, dstData, batchSize, c, h, w, stride, 1, 1, stream); + return 0; +} +#endif size_t UpsampleRT::getSerializationSize() const NOEXCEPT { return 4*sizeof(int); @@ -97,9 +109,8 @@ IPluginV2 *UpsampleRT::clone() const NOEXCEPT { return p; } - UpsampleRTPluginCreator::UpsampleRTPluginCreator() { - mPluginAttributes.emplace_back(PluginField("stride",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); mFC.fields = mPluginAttributes.data(); } diff --git a/src/pluginsRT/YoloRT.cpp b/src/pluginsRT/YoloRT.cpp index 7f615b4..c191afd 100644 --- a/src/pluginsRT/YoloRT.cpp +++ b/src/pluginsRT/YoloRT.cpp @@ -85,6 +85,7 @@ size_t YoloRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } +#if NV_TENSORRT_MAJOR > 7 int YoloRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { dnnType *srcData = (dnnType *) reinterpret_cast(inputs[0]); @@ -115,6 +116,38 @@ int YoloRT::enqueue(int batchSize, const void *const *inputs, void *const *outpu //std::cout<<"YOLO END\n"; return 0; } +#elif NV_TENSORRT_MAJOR == 7 +int32_t YoloRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) { + dnnType *srcData = (dnnType *) reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + + checkCuda(cudaMemcpyAsync(dstData, srcData, batchSize * c * h * w * sizeof(dnnType), cudaMemcpyDeviceToDevice, + stream)); + + + for (int b = 0; b < batchSize; ++b) { + for (int n = 0; n < n_masks; ++n) { + int index = entry_index(b, n * w * h, 0); + if (new_coords == 1) { + if (this->scaleXY != 1) + scalAdd(dstData + index, 2 * w * h, this->scaleXY, -0.5 * (this->scaleXY - 1), 1); + } else { + activationLOGISTICForward(srcData + index, dstData + index, 2 * w * h, stream); //x,y + + if (this->scaleXY != 1) + scalAdd(dstData + index, 2 * w * h, this->scaleXY, -0.5 * (this->scaleXY - 1), 1); + + index = entry_index(b, n * w * h, 4); + activationLOGISTICForward(srcData + index, dstData + index, (1 + classes) * w * h, stream); + } + } + } + + //std::cout<<"YOLO END\n"; + return 0; +} +#endif + size_t YoloRT::getSerializationSize() const NOEXCEPT { return 8 * sizeof(int) + 2 * sizeof(float) + n_masks * sizeof(dnnType) + num * n_masks * 2 * sizeof(dnnType) + @@ -182,16 +215,8 @@ IPluginV2 *YoloRT::clone() const NOEXCEPT { return p; } - YoloRTPluginCreator::YoloRTPluginCreator() { - mPluginAttributes.emplace_back(PluginField("classes",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("num",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("yolo",nullptr,PluginFieldType::kUNKNOWN,1)); - mPluginAttributes.emplace_back(PluginField("numMasks",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("scaleXY",nullptr,PluginFieldType::kFLOAT32,1)); - mPluginAttributes.emplace_back(PluginField("nmsThresh",nullptr,PluginFieldType::kFLOAT32,1)); - mPluginAttributes.emplace_back(PluginField("nmsKind",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("newCoords",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.clear(); mFC.nbFields = mPluginAttributes.size(); mFC.fields = mPluginAttributes.data(); }