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..a8138a1 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -7,10 +7,11 @@ #include "Layer.h" #include "NvInfer.h" #include -#include +#include #include #include #include +#include #include #include #include @@ -30,85 +31,85 @@ 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/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..549e5a0 --- /dev/null +++ b/include/tkDNN/pluginsRT/ActivationSwishRT.h @@ -0,0 +1,82 @@ +#include +#include "../kernels.h" +#include +#include + +namespace nvinfer1 { + class ActivationSwishRT : public IPluginV2 { + + public: + ActivationSwishRT() ; + + ~ActivationSwishRT() ; + + ActivationSwishRT(const void *data, size_t length) ; + + + int getNbOutputs() const NOEXCEPT override ; + + 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 initialize() NOEXCEPT override ; + + 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 ; + + void serialize(void *buffer) const NOEXCEPT override ; + + const char *getPluginType() const NOEXCEPT override ; + + const char *getPluginVersion() const NOEXCEPT override ; + + void destroy() NOEXCEPT override { delete this; } + + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ; + + const char *getPluginNamespace() const NOEXCEPT override ; + + void setPluginNamespace(const char *plguinNamespace) NOEXCEPT override ; + + IPluginV2 *clone() const NOEXCEPT override ; + + 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/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..9b14970 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); @@ -646,8 +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 == CUDNN_ACTIVATION_ELU || l->act_mode == ACTIVATION_ELU) { + IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kELU); + } + else if(l->act_mode == ACTIVATION_SWISH) { + IPluginV2 *plugin = new ActivationSwishRT(); + ILayer *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; } @@ -1025,18 +1029,17 @@ bool NetworkRT::deserialize(const char *filename) { } #if NV_TENSORRT_MAJOR > 7 -void NetworkRT::destroy() { - delete contextRT; - if(builderActive) { - delete engineRT; - delete builderRT; - } -} + 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/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); +} 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; +} + 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; +}