From 000af24d29e1d406f55f3d5f24aa4b01014bfeb4 Mon Sep 17 00:00:00 2001 From: "Cuellar, Adam T" Date: Tue, 27 Jul 2021 13:09:43 +0000 Subject: [PATCH 1/8] Add swish --- include/tkDNN/Layer.h | 7 ++- include/tkDNN/NetworkRT.h | 1 + include/tkDNN/kernels.h | 1 + include/tkDNN/pluginsRT/ActivationSwishRT.h | 61 +++++++++++++++++++++ src/Activation.cpp | 4 ++ src/DarknetParser.cpp | 1 + src/NetworkRT.cpp | 14 ++++- src/kernels/activation_swish.cu | 23 ++++++++ 8 files changed, 110 insertions(+), 2 deletions(-) create mode 100644 include/tkDNN/pluginsRT/ActivationSwishRT.h create mode 100644 src/kernels/activation_swish.cu diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index d1234a5..a8bcecc 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -19,6 +19,7 @@ enum layerType_t { LAYER_ACTIVATION_CRELU, LAYER_ACTIVATION_LEAKY, LAYER_ACTIVATION_MISH, + LAYER_ACTIVATION_SWISH, LAYER_ACTIVATION_LOGISTIC, LAYER_FLATTEN, LAYER_RESHAPE, @@ -74,6 +75,7 @@ public: case LAYER_ACTIVATION_CRELU: return "ActivationCReLU"; case LAYER_ACTIVATION_LEAKY: return "ActivationLeaky"; case LAYER_ACTIVATION_MISH: return "ActivationMish"; + case LAYER_ACTIVATION_SWISH: return "ActivationSwish"; case LAYER_ACTIVATION_LOGISTIC: return "ActivationLogistic"; case LAYER_FLATTEN: return "Flatten"; case LAYER_RESHAPE: return "Reshape"; @@ -221,7 +223,8 @@ typedef enum { ACTIVATION_ELU = 100, ACTIVATION_LEAKY = 101, ACTIVATION_MISH = 102, - ACTIVATION_LOGISTIC = 103 + ACTIVATION_LOGISTIC = 103, + ACTIVATION_SWISH = 104 } tkdnnActivationMode_t; /** @@ -243,6 +246,8 @@ public: return LAYER_ACTIVATION_LEAKY; else if (act_mode == ACTIVATION_MISH) return LAYER_ACTIVATION_MISH; + else if (act_mode == ACTIVATION_SWISH) + return LAYER_ACTIVATION_SWISH; else if (act_mode == ACTIVATION_LOGISTIC) return LAYER_ACTIVATION_LOGISTIC; else diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 9892a24..e5bdd39 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -28,6 +28,7 @@ using namespace nvinfer1; #include "pluginsRT/ActivationLogisticRT.h" #include "pluginsRT/ActivationReLUCeilingRT.h" #include "pluginsRT/ActivationMishRT.h" +#include "pluginsRT/ActivationSwishRT.h" #include "pluginsRT/ReorgRT.h" #include "pluginsRT/RegionRT.h" #include "pluginsRT/RouteRT.h" diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index d809129..52f8bb3 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -9,6 +9,7 @@ void activationReLUCeilingForward(dnnType *srcData, dnnType *dstData, int size, void activationLOGISTICForward(dnnType *srcData, dnnType *dstData, int size, cudaStream_t stream = cudaStream_t(0)); void activationSIGMOIDForward(dnnType *srcData, dnnType *dstData, int size, cudaStream_t stream = cudaStream_t(0)); void activationMishForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream= cudaStream_t(0)); +void activationSwishForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream= cudaStream_t(0)); void fill(dnnType *data, int size, dnnType val, cudaStream_t stream = cudaStream_t(0)); diff --git a/include/tkDNN/pluginsRT/ActivationSwishRT.h b/include/tkDNN/pluginsRT/ActivationSwishRT.h new file mode 100644 index 0000000..123ff1f --- /dev/null +++ b/include/tkDNN/pluginsRT/ActivationSwishRT.h @@ -0,0 +1,61 @@ +#include +#include "../kernels.h" + +class ActivationSwishRT : public IPlugin { + +public: + ActivationSwishRT() { + + + } + + ~ActivationSwishRT(){ + + } + + int getNbOutputs() const override { + return 1; + } + + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return inputs[0]; + } + + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + size = 1; + for(int i=0; i(inputs[0]), + reinterpret_cast(outputs[0]), batchSize*size, stream); + return 0; + } + + + virtual size_t getSerializationSize() override { + return 1*sizeof(int); + } + + virtual void serialize(void* buffer) override { + char *buf = reinterpret_cast(buffer),*a=buf; + tk::dnn::writeBUF(buf, size); + assert(buf == a + getSerializationSize()); + } + + int size; +}; diff --git a/src/Activation.cpp b/src/Activation.cpp index 0b113a7..62c26bd 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -52,6 +52,10 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) { else if(act_mode == ACTIVATION_MISH) { activationMishForward(srcData, dstData, dim.tot()); + } + else if(act_mode == ACTIVATION_SWISH) { + activationSwishForward(srcData, dstData, dim.tot()); + } else if(act_mode == ACTIVATION_LOGISTIC) { activationLOGISTICForward(srcData, dstData, dim.tot()); diff --git a/src/DarknetParser.cpp b/src/DarknetParser.cpp index 69b6b29..247289c 100644 --- a/src/DarknetParser.cpp +++ b/src/DarknetParser.cpp @@ -187,6 +187,7 @@ namespace tk { namespace dnn { if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU); else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY; else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH; + else if(f.activation == "swish") act = tk::dnn::ACTIVATION_SWISH; else if(f.activation == "logistic") act = tk::dnn::ACTIVATION_LOGISTIC; else { FatalError("activation not supported: " + f.activation); } netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act); diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 6ac7235..810c117 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -227,7 +227,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) { return convert_layer(input, (Conv2d*) l); if(type == LAYER_POOLING) return convert_layer(input, (Pooling*) l); - if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_LOGISTIC) + if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_SWISH || type == LAYER_ACTIVATION_LOGISTIC) return convert_layer(input, (Activation*) l); if(type == LAYER_SOFTMAX) return convert_layer(input, (Softmax*) l); @@ -424,6 +424,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { checkNULL(lRT); return lRT; } + else if(l->act_mode == ACTIVATION_SWISH) { + IPlugin *plugin = new ActivationSwishRT(); + IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + checkNULL(lRT); + return lRT; + } else if(l->act_mode == ACTIVATION_LOGISTIC) { IPlugin *plugin = new ActivationLogisticRT(); IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); @@ -674,6 +680,12 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa assert(buf == bufCheck + serialLength); return a; } + if(name.find("ActivationSwish") == 0) { + ActivationSwishRT *a = new ActivationSwishRT(); + a->size = readBUF(buf); + assert(buf == bufCheck + serialLength); + return a; + } if(name.find("ActivationLogistic") == 0) { ActivationLogisticRT *a = new ActivationLogisticRT(); a->size = readBUF(buf); diff --git a/src/kernels/activation_swish.cu b/src/kernels/activation_swish.cu new file mode 100644 index 0000000..7018db6 --- /dev/null +++ b/src/kernels/activation_swish.cu @@ -0,0 +1,23 @@ +#include "kernels.h" + +__global__ +void activation_swish(dnnType *input, dnnType *output, int size) { + + int i = blockDim.x*blockIdx.x + threadIdx.x; + + if(i>>(srcData, dstData, size); +} -- 2.52.0 From 2e8ffb9d1dcc1225af5422fb37299c2c1b75a640 Mon Sep 17 00:00:00 2001 From: AdamCuellar Date: Wed, 15 Sep 2021 08:24:13 -0400 Subject: [PATCH 2/8] Add yolov4-csp-swish --- tests/darknet/cfg/yolov4-csp-swish.cfg | 1355 ++++++++++++++++++++++++ tests/darknet/yolov4-csp-swish.cpp | 40 + 2 files changed, 1395 insertions(+) create mode 100644 tests/darknet/cfg/yolov4-csp-swish.cfg create mode 100644 tests/darknet/yolov4-csp-swish.cpp diff --git a/tests/darknet/cfg/yolov4-csp-swish.cfg b/tests/darknet/cfg/yolov4-csp-swish.cfg new file mode 100644 index 0000000..2aab444 --- /dev/null +++ b/tests/darknet/cfg/yolov4-csp-swish.cfg @@ -0,0 +1,1355 @@ +[net] +# Testing +#batch=1 +#subdivisions=1 +# Training +batch=64 +subdivisions=8 +width=640 +height=640 +channels=3 +momentum=0.949 +decay=0.0005 +angle=0 +saturation = 1.5 +exposure = 1.5 +hue=.1 + +learning_rate=0.001 +burn_in=1000 +max_batches = 500500 +policy=steps +steps=400000,450000 +scales=.1,.1 + +mosaic=1 + +letter_box=1 + +ema_alpha=0.9998 + +#optimized_memory=1 + + +# ============ Backbone ============ # + +# Stem + +# 0 +[convolutional] +batch_normalize=1 +filters=32 +size=3 +stride=1 +pad=1 +activation=swish + +# P1 + +# Downsample + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=2 +pad=1 +activation=swish + +# Residual Block + +[convolutional] +batch_normalize=1 +filters=32 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=swish + +# 4 (previous+1+3k) +[shortcut] +from=-3 +activation=linear + +# P2 + +# Downsample + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=2 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=swish + +# Residual Block + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +# Transition first + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=swish + +# Merge [-1, -(3k+4)] + +[route] +layers = -1,-10 + +# Transition last + +# 17 (previous+7+3k) +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +# P3 + +# Downsample + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=2 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +# Residual Block + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +# Transition first + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +# Merge [-1 -(4+3k)] + +[route] +layers = -1,-28 + +# Transition last + +# 48 (previous+7+3k) +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +# P4 + +# Downsample + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=2 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +# Residual Block + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +# Transition first + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +# Merge [-1 -(3k+4)] + +[route] +layers = -1,-28 + +# Transition last + +# 79 (previous+7+3k) +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +# P5 + +# Downsample + +[convolutional] +batch_normalize=1 +filters=1024 +size=3 +stride=2 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +# Residual Block + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +# Transition first + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +# Merge [-1 -(3k+4)] + +[route] +layers = -1,-16 + +# Transition last + +# 98 (previous+7+3k) +[convolutional] +batch_normalize=1 +filters=1024 +size=1 +stride=1 +pad=1 +activation=swish + +# ============ End of Backbone ============ # + +# ============ Neck ============ # + +# CSPSPP + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=swish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +### SPP ### +[maxpool] +stride=1 +size=5 + +[route] +layers=-2 + +[maxpool] +stride=1 +size=9 + +[route] +layers=-4 + +[maxpool] +stride=1 +size=13 + +[route] +layers=-1,-3,-5,-6 +### End SPP ### + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=swish + +[route] +layers = -1, -13 + +# 113 (previous+6+5+2k) +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +# End of CSPSPP + + +# FPN-4 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[upsample] +stride=2 + +[route] +layers = 79 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -1, -3 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +# Plain Block + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=swish + +# Merge [-1, -(2k+2)] + +[route] +layers = -1, -6 + +# Transition last + +# 127 (previous+6+4+2k) +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + + +# FPN-3 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[upsample] +stride=2 + +[route] +layers = 48 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -1, -3 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +# Plain Block + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=128 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=128 +activation=swish + +# Merge [-1, -(2k+2)] + +[route] +layers = -1, -6 + +# Transition last + +# 141 (previous+6+4+2k) +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + + +# PAN-4 + +[convolutional] +batch_normalize=1 +size=3 +stride=2 +pad=1 +filters=256 +activation=swish + +[route] +layers = -1, 127 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +# Plain Block + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=swish + +[route] +layers = -1,-6 + +# Transition last + +# 152 (previous+3+4+2k) +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + + +# PAN-5 + +[convolutional] +batch_normalize=1 +size=3 +stride=2 +pad=1 +filters=512 +activation=swish + +[route] +layers = -1, 113 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +# Plain Block + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=swish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=swish + +[route] +layers = -1,-6 + +# Transition last + +# 163 (previous+3+4+2k) +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish +stopbackward=900 + +# ============ End of Neck ============ # + +# ============ Head ============ # + +# YOLO-3 + +[route] +layers = 141 + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=swish + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=logistic + +[yolo] +mask = 0,1,2 +anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 459, 401 +classes=80 +num=9 +jitter=.1 +scale_x_y = 2.0 +objectness_smooth=1 +ignore_thresh = .7 +truth_thresh = 1 +#random=1 +resize=1.5 +#iou_thresh=0.2 +iou_normalizer=0.05 +cls_normalizer=0.5 +obj_normalizer=0.4 +iou_loss=ciou +nms_kind=diounms +beta_nms=0.6 +new_coords=1 +max_delta=2 + + +# YOLO-4 + +[route] +layers = 152 + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=swish + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=logistic + +[yolo] +mask = 3,4,5 +anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 459, 401 +classes=80 +num=9 +jitter=.1 +scale_x_y = 2.0 +objectness_smooth=1 +ignore_thresh = .7 +truth_thresh = 1 +#random=1 +resize=1.5 +#iou_thresh=0.2 +iou_normalizer=0.05 +cls_normalizer=0.5 +obj_normalizer=0.4 +iou_loss=ciou +nms_kind=diounms +beta_nms=0.6 +new_coords=1 +max_delta=2 + + +# YOLO-5 + +[route] +layers = 163 + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=1024 +activation=swish + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=logistic + +[yolo] +mask = 6,7,8 +anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 459, 401 +classes=80 +num=9 +jitter=.1 +scale_x_y = 2.0 +objectness_smooth=1 +ignore_thresh = .7 +truth_thresh = 1 +#random=1 +resize=1.5 +#iou_thresh=0.2 +iou_normalizer=0.05 +cls_normalizer=0.5 +obj_normalizer=0.4 +iou_loss=ciou +nms_kind=diounms +beta_nms=0.6 +new_coords=1 +max_delta=2 diff --git a/tests/darknet/yolov4-csp-swish.cpp b/tests/darknet/yolov4-csp-swish.cpp new file mode 100644 index 0000000..b097ba2 --- /dev/null +++ b/tests/darknet/yolov4-csp-swish.cpp @@ -0,0 +1,40 @@ +// +// Created by Adam T. Cuellar on 9/15/21. +// + +#include +#include +#include "tkdnn.h" +#include "test.h" +#include "DarknetParser.h" + +int main() { + std::string bin_path = "yolov4-csp-swish"; + std::vector input_bins = { + bin_path + "/layers/input.bin" + }; + std::vector output_bins = { + bin_path + "/debug/layer167_out.bin", + bin_path + "/debug/layer171_out.bin", + bin_path + "/debug/layer175_out.bin" + }; + std::string wgs_path = bin_path + "/layers"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolov4-csp-swish.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; + downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/5MFjtNtgbDGdJEo/download"); + + + + // parse darknet network + tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path); + net->print(); + + //convert network to tensorRT + tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str())); + + int ret = testInference(input_bins, output_bins, net, netRT); + net->releaseLayers(); + delete net; + delete netRT; + return ret; +} -- 2.52.0 From bacabb4b42fcbc86f29a6250d09cbe23f5362ced Mon Sep 17 00:00:00 2001 From: "Cuellar, Adam T" Date: Mon, 4 Oct 2021 19:41:35 +0000 Subject: [PATCH 3/8] Update cmake --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index b366416..0a390c6 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -47,7 +47,7 @@ target_link_libraries(kernels ${CUDA_CUBLAS_LIBRARIES}) #------------------------------------------------------------------------------- # External Libraries #------------------------------------------------------------------------------- -find_package(Eigen3 REQUIRED) +find_package(Eigen3 3.4 REQUIRED) message("Eigen DIR: " ${EIGEN3_INCLUDE_DIR}) include_directories(${EIGEN3_INCLUDE_DIR}) -- 2.52.0 From 69ff4f37f4c05de9f36b35e6f34e9fe67e95520a Mon Sep 17 00:00:00 2001 From: "Cuellar, Adam T" Date: Tue, 27 Jul 2021 13:09:43 +0000 Subject: [PATCH 4/8] Add swish --- include/tkDNN/Layer.h | 7 ++- include/tkDNN/NetworkRT.h | 41 ++++++++++++++ include/tkDNN/kernels.h | 1 + include/tkDNN/pluginsRT/ActivationSwishRT.h | 61 +++++++++++++++++++++ src/Activation.cpp | 4 ++ src/DarknetParser.cpp | 1 + src/NetworkRT.cpp | 41 +++++++++++++- src/kernels/activation_swish.cu | 23 ++++++++ 8 files changed, 177 insertions(+), 2 deletions(-) create mode 100644 include/tkDNN/pluginsRT/ActivationSwishRT.h create mode 100644 src/kernels/activation_swish.cu diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index daa27e5..236cafd 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -19,6 +19,7 @@ enum layerType_t { LAYER_ACTIVATION_CRELU, LAYER_ACTIVATION_LEAKY, LAYER_ACTIVATION_MISH, + LAYER_ACTIVATION_SWISH, LAYER_ACTIVATION_LOGISTIC, LAYER_FLATTEN, LAYER_RESHAPE, @@ -75,6 +76,7 @@ public: case LAYER_ACTIVATION_CRELU: return "ActivationCReLU"; case LAYER_ACTIVATION_LEAKY: return "ActivationLeaky"; case LAYER_ACTIVATION_MISH: return "ActivationMish"; + case LAYER_ACTIVATION_SWISH: return "ActivationSwish"; case LAYER_ACTIVATION_LOGISTIC: return "ActivationLogistic"; case LAYER_FLATTEN: return "Flatten"; case LAYER_RESHAPE: return "Reshape"; @@ -223,7 +225,8 @@ typedef enum { ACTIVATION_ELU = 100, ACTIVATION_LEAKY = 101, ACTIVATION_MISH = 102, - ACTIVATION_LOGISTIC = 103 + ACTIVATION_LOGISTIC = 103, + ACTIVATION_SWISH = 104 } tkdnnActivationMode_t; /** @@ -245,6 +248,8 @@ public: return LAYER_ACTIVATION_LEAKY; else if (act_mode == ACTIVATION_MISH) return LAYER_ACTIVATION_MISH; + else if (act_mode == ACTIVATION_SWISH) + return LAYER_ACTIVATION_SWISH; else if (act_mode == ACTIVATION_LOGISTIC) return LAYER_ACTIVATION_LOGISTIC; else diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index a422134..5e8d980 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -26,6 +26,47 @@ #include #include +namespace tk { namespace dnn { + +template void writeBUF(char*& buffer, const T& val) +{ + *reinterpret_cast(buffer) = val; + buffer += sizeof(T); +} + +template T readBUF(const char*& buffer) +{ + T val = *reinterpret_cast(buffer); + buffer += sizeof(T); + return val; +} + +using namespace nvinfer1; +#include "pluginsRT/ActivationLeakyRT.h" +#include "pluginsRT/ActivationLogisticRT.h" +#include "pluginsRT/ActivationReLUCeilingRT.h" +#include "pluginsRT/ActivationMishRT.h" +#include "pluginsRT/ActivationSwishRT.h" +#include "pluginsRT/ReorgRT.h" +#include "pluginsRT/RegionRT.h" +#include "pluginsRT/RouteRT.h" +#include "pluginsRT/ShortcutRT.h" +#include "pluginsRT/YoloRT.h" +#include "pluginsRT/UpsampleRT.h" +#include "pluginsRT/ResizeLayerRT.h" +#include "pluginsRT/DeformableConvRT.h" +#include "pluginsRT/FlattenConcatRT.h" +#include "pluginsRT/ReshapeRT.h" +#include "pluginsRT/MaxPoolingFixedSizeRT.h" + +class PluginFactory : IPluginFactory +{ +public: + YoloRT *yolos[16]; + int n_yolos; + + virtual IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength); +}; namespace tk { namespace dnn { diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index 4d5474b..16bb033 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -9,6 +9,7 @@ void activationReLUCeilingForward(dnnType *srcData, dnnType *dstData, int size, void activationLOGISTICForward(dnnType *srcData, dnnType *dstData, int size, cudaStream_t stream = cudaStream_t(0)); void activationSIGMOIDForward(dnnType *srcData, dnnType *dstData, int size, cudaStream_t stream = cudaStream_t(0)); void activationMishForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream= cudaStream_t(0)); +void activationSwishForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream= cudaStream_t(0)); void fill(dnnType *data, int size, dnnType val, cudaStream_t stream = cudaStream_t(0)); diff --git a/include/tkDNN/pluginsRT/ActivationSwishRT.h b/include/tkDNN/pluginsRT/ActivationSwishRT.h new file mode 100644 index 0000000..123ff1f --- /dev/null +++ b/include/tkDNN/pluginsRT/ActivationSwishRT.h @@ -0,0 +1,61 @@ +#include +#include "../kernels.h" + +class ActivationSwishRT : public IPlugin { + +public: + ActivationSwishRT() { + + + } + + ~ActivationSwishRT(){ + + } + + int getNbOutputs() const override { + return 1; + } + + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return inputs[0]; + } + + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + size = 1; + for(int i=0; i(inputs[0]), + reinterpret_cast(outputs[0]), batchSize*size, stream); + return 0; + } + + + virtual size_t getSerializationSize() override { + return 1*sizeof(int); + } + + virtual void serialize(void* buffer) override { + char *buf = reinterpret_cast(buffer),*a=buf; + tk::dnn::writeBUF(buf, size); + assert(buf == a + getSerializationSize()); + } + + int size; +}; diff --git a/src/Activation.cpp b/src/Activation.cpp index 947b019..49fde03 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -52,6 +52,10 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) { else if(act_mode == ACTIVATION_MISH) { activationMishForward(srcData, dstData, dim.tot()); + } + else if(act_mode == ACTIVATION_SWISH) { + activationSwishForward(srcData, dstData, dim.tot()); + } else if(act_mode == ACTIVATION_LOGISTIC) { activationLOGISTICForward(srcData, dstData, dim.tot()); diff --git a/src/DarknetParser.cpp b/src/DarknetParser.cpp index 3333afd..7006246 100644 --- a/src/DarknetParser.cpp +++ b/src/DarknetParser.cpp @@ -197,6 +197,7 @@ namespace tk { namespace dnn { if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU); else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY; else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH; + else if(f.activation == "swish") act = tk::dnn::ACTIVATION_SWISH; else if(f.activation == "logistic") act = tk::dnn::ACTIVATION_LOGISTIC; else { FatalError("activation not supported: " + f.activation); } netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act); diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 26489bf..9617775 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -254,7 +254,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) { return convert_layer(input, (Conv2d*) l); if(type == LAYER_POOLING) return convert_layer(input, (Pooling*) l); - if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_LOGISTIC) + if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_SWISH || type == LAYER_ACTIVATION_LOGISTIC) return convert_layer(input, (Activation*) l); if(type == LAYER_SOFTMAX) return convert_layer(input, (Softmax*) l); @@ -648,6 +648,15 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { } else if(l->act_mode == CUDNN_ACTIVATION_ELU || l->act_mode == ACTIVATION_ELU){ IActivationLayer *lRT = networkRT->addActivation(*input,ActivationType::kELU); + else if(l->act_mode == ACTIVATION_SWISH) { + IPlugin *plugin = new ActivationSwishRT(); + IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + checkNULL(lRT); + return lRT; + } + else if(l->act_mode == ACTIVATION_LOGISTIC) { + IPlugin *plugin = new ActivationLogisticRT(); + IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); checkNULL(lRT); return lRT; } @@ -1030,6 +1039,36 @@ void NetworkRT::destroy() { if(builderActive) { delete engineRT; delete builderRT; + + +IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) { + const char * buf = reinterpret_cast(serialData),*bufCheck = buf; + + std::string name(layerName); + //std::cout<(buf)); + a->size = readBUF(buf); + assert(buf == bufCheck + serialLength); + return a; + } + if(name.find("ActivationMish") == 0) { + ActivationMishRT *a = new ActivationMishRT(); + a->size = readBUF(buf); + assert(buf == bufCheck + serialLength); + return a; + } + if(name.find("ActivationSwish") == 0) { + ActivationSwishRT *a = new ActivationSwishRT(); + a->size = readBUF(buf); + assert(buf == bufCheck + serialLength); + return a; + } + if(name.find("ActivationLogistic") == 0) { + ActivationLogisticRT *a = new ActivationLogisticRT(); + a->size = readBUF(buf); + return a; } } #elif NV_TENSORRT_MAJOR <=7 diff --git a/src/kernels/activation_swish.cu b/src/kernels/activation_swish.cu new file mode 100644 index 0000000..7018db6 --- /dev/null +++ b/src/kernels/activation_swish.cu @@ -0,0 +1,23 @@ +#include "kernels.h" + +__global__ +void activation_swish(dnnType *input, dnnType *output, int size) { + + int i = blockDim.x*blockIdx.x + threadIdx.x; + + if(i>>(srcData, dstData, size); +} -- 2.52.0 From c8bb81b23687a367fedf5b3531b1e4a00c5ca772 Mon Sep 17 00:00:00 2001 From: "Cuellar, Adam T" Date: Mon, 4 Oct 2021 19:41:35 +0000 Subject: [PATCH 5/8] Update cmake --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index c29e987..75c785b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -102,7 +102,7 @@ target_link_libraries(kernels ${CUDA_CUBLAS_LIBRARIES} ${CUDA_LIBRARIES} ${CUDNN #------------------------------------------------------------------------------- # External Libraries #------------------------------------------------------------------------------- -find_package(Eigen3 REQUIRED) +find_package(Eigen3 3.4 REQUIRED) message("Eigen DIR: " ${EIGEN3_INCLUDE_DIR}) include_directories(${EIGEN3_INCLUDE_DIR}) -- 2.52.0 From a3d607a210799b052733c8ce707cbd41f287887f Mon Sep 17 00:00:00 2001 From: AdamCuellar Date: Wed, 15 Sep 2021 08:24:13 -0400 Subject: [PATCH 6/8] Add yolov4-csp-swish --- tests/darknet/cfg/yolov4-csp-swish.cfg | 1355 ++++++++++++++++++++++++ tests/darknet/yolov4-csp-swish.cpp | 40 + 2 files changed, 1395 insertions(+) create mode 100644 tests/darknet/cfg/yolov4-csp-swish.cfg create mode 100644 tests/darknet/yolov4-csp-swish.cpp diff --git a/tests/darknet/cfg/yolov4-csp-swish.cfg b/tests/darknet/cfg/yolov4-csp-swish.cfg new file mode 100644 index 0000000..2aab444 --- /dev/null +++ b/tests/darknet/cfg/yolov4-csp-swish.cfg @@ -0,0 +1,1355 @@ +[net] +# Testing +#batch=1 +#subdivisions=1 +# Training +batch=64 +subdivisions=8 +width=640 +height=640 +channels=3 +momentum=0.949 +decay=0.0005 +angle=0 +saturation = 1.5 +exposure = 1.5 +hue=.1 + +learning_rate=0.001 +burn_in=1000 +max_batches = 500500 +policy=steps +steps=400000,450000 +scales=.1,.1 + +mosaic=1 + +letter_box=1 + +ema_alpha=0.9998 + +#optimized_memory=1 + + +# ============ Backbone ============ # + +# Stem + +# 0 +[convolutional] +batch_normalize=1 +filters=32 +size=3 +stride=1 +pad=1 +activation=swish + +# P1 + +# Downsample + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=2 +pad=1 +activation=swish + +# Residual Block + +[convolutional] +batch_normalize=1 +filters=32 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=swish + +# 4 (previous+1+3k) +[shortcut] +from=-3 +activation=linear + +# P2 + +# Downsample + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=2 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=swish + +# Residual Block + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +# Transition first + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=swish + +# Merge [-1, -(3k+4)] + +[route] +layers = -1,-10 + +# Transition last + +# 17 (previous+7+3k) +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +# P3 + +# Downsample + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=2 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +# Residual Block + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +# Transition first + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +# Merge [-1 -(4+3k)] + +[route] +layers = -1,-28 + +# Transition last + +# 48 (previous+7+3k) +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +# P4 + +# Downsample + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=2 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +# Residual Block + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +# Transition first + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +# Merge [-1 -(3k+4)] + +[route] +layers = -1,-28 + +# Transition last + +# 79 (previous+7+3k) +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +# P5 + +# Downsample + +[convolutional] +batch_normalize=1 +filters=1024 +size=3 +stride=2 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +# Residual Block + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=swish + +[shortcut] +from=-3 +activation=linear + +# Transition first + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +# Merge [-1 -(3k+4)] + +[route] +layers = -1,-16 + +# Transition last + +# 98 (previous+7+3k) +[convolutional] +batch_normalize=1 +filters=1024 +size=1 +stride=1 +pad=1 +activation=swish + +# ============ End of Backbone ============ # + +# ============ Neck ============ # + +# CSPSPP + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=swish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +### SPP ### +[maxpool] +stride=1 +size=5 + +[route] +layers=-2 + +[maxpool] +stride=1 +size=9 + +[route] +layers=-4 + +[maxpool] +stride=1 +size=13 + +[route] +layers=-1,-3,-5,-6 +### End SPP ### + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=swish + +[route] +layers = -1, -13 + +# 113 (previous+6+5+2k) +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +# End of CSPSPP + + +# FPN-4 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[upsample] +stride=2 + +[route] +layers = 79 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -1, -3 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +# Plain Block + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=swish + +# Merge [-1, -(2k+2)] + +[route] +layers = -1, -6 + +# Transition last + +# 127 (previous+6+4+2k) +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + + +# FPN-3 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[upsample] +stride=2 + +[route] +layers = 48 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -1, -3 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +# Plain Block + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=128 +activation=swish + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=128 +activation=swish + +# Merge [-1, -(2k+2)] + +[route] +layers = -1, -6 + +# Transition last + +# 141 (previous+6+4+2k) +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=swish + + +# PAN-4 + +[convolutional] +batch_normalize=1 +size=3 +stride=2 +pad=1 +filters=256 +activation=swish + +[route] +layers = -1, 127 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +# Plain Block + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=swish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=swish + +[route] +layers = -1,-6 + +# Transition last + +# 152 (previous+3+4+2k) +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=swish + + +# PAN-5 + +[convolutional] +batch_normalize=1 +size=3 +stride=2 +pad=1 +filters=512 +activation=swish + +[route] +layers = -1, 113 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +# Split + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[route] +layers = -2 + +# Plain Block + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=swish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=swish + +[route] +layers = -1,-6 + +# Transition last + +# 163 (previous+3+4+2k) +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=swish +stopbackward=900 + +# ============ End of Neck ============ # + +# ============ Head ============ # + +# YOLO-3 + +[route] +layers = 141 + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=swish + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=logistic + +[yolo] +mask = 0,1,2 +anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 459, 401 +classes=80 +num=9 +jitter=.1 +scale_x_y = 2.0 +objectness_smooth=1 +ignore_thresh = .7 +truth_thresh = 1 +#random=1 +resize=1.5 +#iou_thresh=0.2 +iou_normalizer=0.05 +cls_normalizer=0.5 +obj_normalizer=0.4 +iou_loss=ciou +nms_kind=diounms +beta_nms=0.6 +new_coords=1 +max_delta=2 + + +# YOLO-4 + +[route] +layers = 152 + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=swish + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=logistic + +[yolo] +mask = 3,4,5 +anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 459, 401 +classes=80 +num=9 +jitter=.1 +scale_x_y = 2.0 +objectness_smooth=1 +ignore_thresh = .7 +truth_thresh = 1 +#random=1 +resize=1.5 +#iou_thresh=0.2 +iou_normalizer=0.05 +cls_normalizer=0.5 +obj_normalizer=0.4 +iou_loss=ciou +nms_kind=diounms +beta_nms=0.6 +new_coords=1 +max_delta=2 + + +# YOLO-5 + +[route] +layers = 163 + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=1024 +activation=swish + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=logistic + +[yolo] +mask = 6,7,8 +anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 459, 401 +classes=80 +num=9 +jitter=.1 +scale_x_y = 2.0 +objectness_smooth=1 +ignore_thresh = .7 +truth_thresh = 1 +#random=1 +resize=1.5 +#iou_thresh=0.2 +iou_normalizer=0.05 +cls_normalizer=0.5 +obj_normalizer=0.4 +iou_loss=ciou +nms_kind=diounms +beta_nms=0.6 +new_coords=1 +max_delta=2 diff --git a/tests/darknet/yolov4-csp-swish.cpp b/tests/darknet/yolov4-csp-swish.cpp new file mode 100644 index 0000000..b097ba2 --- /dev/null +++ b/tests/darknet/yolov4-csp-swish.cpp @@ -0,0 +1,40 @@ +// +// Created by Adam T. Cuellar on 9/15/21. +// + +#include +#include +#include "tkdnn.h" +#include "test.h" +#include "DarknetParser.h" + +int main() { + std::string bin_path = "yolov4-csp-swish"; + std::vector input_bins = { + bin_path + "/layers/input.bin" + }; + std::vector output_bins = { + bin_path + "/debug/layer167_out.bin", + bin_path + "/debug/layer171_out.bin", + bin_path + "/debug/layer175_out.bin" + }; + std::string wgs_path = bin_path + "/layers"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolov4-csp-swish.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; + downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/5MFjtNtgbDGdJEo/download"); + + + + // parse darknet network + tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path); + net->print(); + + //convert network to tensorRT + tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str())); + + int ret = testInference(input_bins, output_bins, net, netRT); + net->releaseLayers(); + delete net; + delete netRT; + return ret; +} -- 2.52.0 From a37c14d71dfca0833b60c211e51eae1010a983b1 Mon Sep 17 00:00:00 2001 From: Adam Cuellar Date: Mon, 11 Apr 2022 15:57:40 -0400 Subject: [PATCH 7/8] Fix bad merge --- include/tkDNN/NetworkRT.h | 170 ++++++++------------ include/tkDNN/pluginsRT/ActivationSwishRT.h | 101 +++++++----- src/NetworkRT.cpp | 64 ++------ src/pluginsRT/ActivationSwishRT.cpp | 133 +++++++++++++++ 4 files changed, 273 insertions(+), 195 deletions(-) create mode 100644 src/pluginsRT/ActivationSwishRT.cpp diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 5e8d980..b607644 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -11,6 +11,7 @@ #include #include #include +#include #include #include #include @@ -26,130 +27,89 @@ #include #include -namespace tk { namespace dnn { - -template void writeBUF(char*& buffer, const T& val) -{ - *reinterpret_cast(buffer) = val; - buffer += sizeof(T); -} - -template T readBUF(const char*& buffer) -{ - T val = *reinterpret_cast(buffer); - buffer += sizeof(T); - return val; -} - -using namespace nvinfer1; -#include "pluginsRT/ActivationLeakyRT.h" -#include "pluginsRT/ActivationLogisticRT.h" -#include "pluginsRT/ActivationReLUCeilingRT.h" -#include "pluginsRT/ActivationMishRT.h" -#include "pluginsRT/ActivationSwishRT.h" -#include "pluginsRT/ReorgRT.h" -#include "pluginsRT/RegionRT.h" -#include "pluginsRT/RouteRT.h" -#include "pluginsRT/ShortcutRT.h" -#include "pluginsRT/YoloRT.h" -#include "pluginsRT/UpsampleRT.h" -#include "pluginsRT/ResizeLayerRT.h" -#include "pluginsRT/DeformableConvRT.h" -#include "pluginsRT/FlattenConcatRT.h" -#include "pluginsRT/ReshapeRT.h" -#include "pluginsRT/MaxPoolingFixedSizeRT.h" - -class PluginFactory : IPluginFactory -{ -public: - YoloRT *yolos[16]; - int n_yolos; - - virtual IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength); -}; namespace tk { namespace dnn { -class NetworkRT { + class NetworkRT { -public: - nvinfer1::DataType dtRT; - nvinfer1::IBuilder *builderRT; - nvinfer1::IRuntime *runtimeRT; - nvinfer1::INetworkDefinition *networkRT; -#if NV_TENSORRT_MAJOR >= 6 - nvinfer1::IBuilderConfig *configRT; + public: + nvinfer1::DataType dtRT; + nvinfer1::IBuilder *builderRT; + nvinfer1::IRuntime *runtimeRT; + nvinfer1::INetworkDefinition *networkRT; +#if NV_TENSORRT_MAJOR >= 6 + nvinfer1::IBuilderConfig *configRT; #endif - - nvinfer1::ICudaEngine *engineRT; - nvinfer1::IExecutionContext *contextRT; - const static int MAX_BUFFERS_RT = 10; - void* buffersRT[MAX_BUFFERS_RT]; - dataDim_t buffersDIM[MAX_BUFFERS_RT]; - int buf_input_idx, buf_output_idx; - bool builderActive = false; - dataDim_t input_dim, output_dim; - dnnType *output; - cudaStream_t stream; + nvinfer1::ICudaEngine *engineRT; + nvinfer1::IExecutionContext *contextRT; - std::vector yolo_plugins; // yolo layers in network + const static int MAX_BUFFERS_RT = 10; + void* buffersRT[MAX_BUFFERS_RT]; + dataDim_t buffersDIM[MAX_BUFFERS_RT]; + int buf_input_idx, buf_output_idx; + bool builderActive = false; + dataDim_t input_dim, output_dim; + dnnType *output; + cudaStream_t stream; - NetworkRT(Network *net, const char *name); - virtual ~NetworkRT(); + std::vector yolo_plugins; // yolo layers in network - int getMaxBatchSize() { - if(engineRT != nullptr) - return engineRT->getMaxBatchSize(); - else - return 0; - } + NetworkRT(Network *net, const char *name); + virtual ~NetworkRT(); - int getBuffersN() { - if(engineRT != nullptr) - return engineRT->getNbBindings(); - else - return 0; - } + int getMaxBatchSize() { + if(engineRT != nullptr) + return engineRT->getMaxBatchSize(); + else + return 0; + } - /** - Do inference - */ - dnnType* infer(dataDim_t &dim, dnnType* data); - void enqueue(int batchSize = 1); + int getBuffersN() { + if(engineRT != nullptr) + return engineRT->getNbBindings(); + else + return 0; + } - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Layer *l); - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Conv2d *l); - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Activation *l); - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Dense *l); - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Pooling *l); - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Softmax *l); - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Route *l); - nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Flatten *l); - nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Reshape *l); - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Resize *l); - nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Reorg *l); - nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Region *l); - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Shortcut *l); - nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Yolo *l); - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Upsample *l); - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, DeformConv2d *l); - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input,Padding *l); - nvinfer1::ILayer* convert_layer(nvinfer1::ITensor* input,MulAdd *l); + /** + Do inference + */ + dnnType* infer(dataDim_t &dim, dnnType* data); + void enqueue(int batchSize = 1); + + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Layer *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Conv2d *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Activation *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Dense *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Pooling *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Softmax *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Route *l); + nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Flatten *l); + nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Reshape *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Resize *l); + nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Reorg *l); + nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Region *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Shortcut *l); + nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Yolo *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Upsample *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, DeformConv2d *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input,Padding *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor* input,MulAdd *l); #if NV_TENSORRT_MAJOR > 5 && NV_TENSORRT_MAJOR < 8 - bool serialize(const char *filename); + bool serialize(const char *filename); #else - bool serialize(const char *filename,nvinfer1::IHostMemory *ptr); + bool serialize(const char *filename,nvinfer1::IHostMemory *ptr); #endif - bool deserialize(const char *filename); - void destroy(); + bool deserialize(const char *filename); + void destroy(); -}; + }; -}} -#endif //NETWORKRT_H + }} +#endif //NETWORKRT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ActivationSwishRT.h b/include/tkDNN/pluginsRT/ActivationSwishRT.h index 123ff1f..549e5a0 100644 --- a/include/tkDNN/pluginsRT/ActivationSwishRT.h +++ b/include/tkDNN/pluginsRT/ActivationSwishRT.h @@ -1,61 +1,82 @@ #include #include "../kernels.h" +#include +#include -class ActivationSwishRT : public IPlugin { +namespace nvinfer1 { + class ActivationSwishRT : public IPluginV2 { -public: - ActivationSwishRT() { + public: + ActivationSwishRT() ; + + ~ActivationSwishRT() ; + + ActivationSwishRT(const void *data, size_t length) ; - } + int getNbOutputs() const NOEXCEPT override ; - ~ActivationSwishRT(){ + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ; - } + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override ; - int getNbOutputs() const override { - return 1; - } + int initialize() NOEXCEPT override ; - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { - return inputs[0]; - } + void terminate() NOEXCEPT override ; - void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { - size = 1; - for(int i=0; i 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 - int initialize() override { + size_t getSerializationSize() const NOEXCEPT override ; - return 0; - } + void serialize(void *buffer) const NOEXCEPT override ; - virtual void terminate() override { - } + const char *getPluginType() const NOEXCEPT override ; - virtual size_t getWorkspaceSize(int maxBatchSize) const override { - return 0; - } + const char *getPluginVersion() const NOEXCEPT override ; - virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { + void destroy() NOEXCEPT override { delete this; } - activationSwishForward((dnnType*)reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), batchSize*size, stream); - return 0; - } + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ; + const char *getPluginNamespace() const NOEXCEPT override ; - virtual size_t getSerializationSize() override { - return 1*sizeof(int); - } + void setPluginNamespace(const char *plguinNamespace) NOEXCEPT override ; - virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, size); - assert(buf == a + getSerializationSize()); - } + IPluginV2 *clone() const NOEXCEPT override ; - int size; -}; + int size; + private: + std::string mPluginNamespace; + }; + + class ActivationSwishRTPluginCreator : public IPluginCreator { + public: + ActivationSwishRTPluginCreator() ; + + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; + const char *getPluginNamespace() const NOEXCEPT override ; + + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ; + + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ; + + const char *getPluginName() const NOEXCEPT override ; + + const char *getPluginVersion() const NOEXCEPT override ; + + const PluginFieldCollection *getFieldNames() NOEXCEPT override ; + + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; + + REGISTER_TENSORRT_PLUGIN(ActivationSwishRTPluginCreator); +}; \ No newline at end of file diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 9617775..9b14970 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -646,17 +646,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { checkNULL(lRT); return lRT; } - else if(l->act_mode == CUDNN_ACTIVATION_ELU || l->act_mode == ACTIVATION_ELU){ - IActivationLayer *lRT = networkRT->addActivation(*input,ActivationType::kELU); - else if(l->act_mode == ACTIVATION_SWISH) { - IPlugin *plugin = new ActivationSwishRT(); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); - checkNULL(lRT); - return lRT; + else if(l->act_mode == CUDNN_ACTIVATION_ELU || l->act_mode == ACTIVATION_ELU) { + IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kELU); } - else if(l->act_mode == ACTIVATION_LOGISTIC) { - IPlugin *plugin = new ActivationLogisticRT(); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + else if(l->act_mode == ACTIVATION_SWISH) { + IPluginV2 *plugin = new ActivationSwishRT(); + ILayer *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; } @@ -1034,48 +1029,17 @@ bool NetworkRT::deserialize(const char *filename) { } #if NV_TENSORRT_MAJOR > 7 -void NetworkRT::destroy() { - delete contextRT; - if(builderActive) { - delete engineRT; - delete builderRT; - - -IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) { - const char * buf = reinterpret_cast(serialData),*bufCheck = buf; - - std::string name(layerName); - //std::cout<(buf)); - a->size = readBUF(buf); - assert(buf == bufCheck + serialLength); - return a; - } - if(name.find("ActivationMish") == 0) { - ActivationMishRT *a = new ActivationMishRT(); - a->size = readBUF(buf); - assert(buf == bufCheck + serialLength); - return a; - } - if(name.find("ActivationSwish") == 0) { - ActivationSwishRT *a = new ActivationSwishRT(); - a->size = readBUF(buf); - assert(buf == bufCheck + serialLength); - return a; - } - if(name.find("ActivationLogistic") == 0) { - ActivationLogisticRT *a = new ActivationLogisticRT(); - a->size = readBUF(buf); - return a; - } -} + void NetworkRT::destroy() { + delete contextRT; + if(builderActive) { + delete engineRT; + delete builderRT; + } + } #elif NV_TENSORRT_MAJOR <=7 -void NetworkRT::destroy() { + void NetworkRT::destroy() { } #endif - -}} +}} \ No newline at end of file diff --git a/src/pluginsRT/ActivationSwishRT.cpp b/src/pluginsRT/ActivationSwishRT.cpp new file mode 100644 index 0000000..3b4b9d0 --- /dev/null +++ b/src/pluginsRT/ActivationSwishRT.cpp @@ -0,0 +1,133 @@ +// +// Created by Adam on 4/11/2022 +// +#include +using namespace nvinfer1; +std::vector ActivationSwishRTPluginCreator::mPluginAttributes; +PluginFieldCollection ActivationSwishRTPluginCreator::mFC{}; + +ActivationSwishRT::ActivationSwishRT() { + +} + +ActivationSwishRT::~ActivationSwishRT() { + +} + +ActivationSwishRT::ActivationSwishRT(const void *data, size_t length) { + const char *buf = reinterpret_cast(data), *bufCheck = buf; + size = readBUF(buf); + assert(buf == bufCheck + length); +} + +int ActivationSwishRT::getNbOutputs() const NOEXCEPT { return 1; } + +Dims ActivationSwishRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { return inputs[0]; } + +void ActivationSwishRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT { +assert(format == PluginFormat::kLINEAR); +size = 1; +for (int i = 0; i < outputDims[0].nbDims; i++) +size *= outputDims[0].d[i]; +} + +int ActivationSwishRT::initialize() NOEXCEPT { return 0; } + +void ActivationSwishRT::terminate() NOEXCEPT {} + +size_t ActivationSwishRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } + +#if NV_TENSORRT_MAJOR > 7 +int ActivationSwishRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + activationSwishForward((dnnType *) reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize * size, stream); + return 0; +} +#elif NV_TENSORRT_MAJOR == 7 +int32_t ActivationSwishRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, + cudaStream_t stream) { + activationSwishForward((dnnType *) reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize * size, stream); + return 0; +} +#endif + +size_t ActivationSwishRT::getSerializationSize() const NOEXCEPT { + return 1 * sizeof(int); +} + +void ActivationSwishRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer), *a = buf; + writeBUF(buf, size); + assert(buf == a + getSerializationSize()); +} + +const char* ActivationSwishRT::getPluginType() const NOEXCEPT { + return "ActivationSwishRT_tkDNN"; +} + +const char *ActivationSwishRT::getPluginVersion() const NOEXCEPT { + return "1"; +} + +bool ActivationSwishRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR); +} + +const char *ActivationSwishRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void ActivationSwishRT::setPluginNamespace(const char *plguinNamespace) NOEXCEPT { + mPluginNamespace = plguinNamespace; +} + +IPluginV2 *ActivationSwishRT::clone() const NOEXCEPT { + auto *p = new ActivationSwishRT(); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + + + +ActivationSwishRTPluginCreator::ActivationSwishRTPluginCreator() { + mPluginAttributes.clear(); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void ActivationSwishRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *ActivationSwishRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *ActivationSwishRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT { + auto *pluginObj = new ActivationSwishRT(serialData, serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *ActivationSwishRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + auto *pluginObj = new ActivationSwishRT(); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *ActivationSwishRTPluginCreator::getPluginName() const NOEXCEPT { + return "ActivationSwishRT_tkDNN"; +} + +const char *ActivationSwishRTPluginCreator::getPluginVersion() const NOEXCEPT{ + return "1"; +} + +const PluginFieldCollection *ActivationSwishRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + -- 2.52.0 From 40dc3f7c2a855e83d18cb7da5d21ed208ec3ac34 Mon Sep 17 00:00:00 2001 From: Adam Cuellar Date: Thu, 25 Aug 2022 13:50:30 -0400 Subject: [PATCH 8/8] Fix eigen version / kernels.h include --- CMakeLists.txt | 2 +- include/tkDNN/NetworkRT.h | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 75c785b..c29e987 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -102,7 +102,7 @@ target_link_libraries(kernels ${CUDA_CUBLAS_LIBRARIES} ${CUDA_LIBRARIES} ${CUDNN #------------------------------------------------------------------------------- # External Libraries #------------------------------------------------------------------------------- -find_package(Eigen3 3.4 REQUIRED) +find_package(Eigen3 REQUIRED) message("Eigen DIR: " ${EIGEN3_INCLUDE_DIR}) include_directories(${EIGEN3_INCLUDE_DIR}) diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index b607644..a8138a1 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -7,7 +7,7 @@ #include "Layer.h" #include "NvInfer.h" #include -#include +#include #include #include #include -- 2.52.0