diff --git a/README.md b/README.md index 6a5281e..4bee544 100644 --- a/README.md +++ b/README.md @@ -86,7 +86,7 @@ Some tricky operations encountered in these models, already solved, but might ha |retinaface output decode| implement a plugin to decode bbox, confidence and landmarks, see retinaface. | |mish| mish activation is implemented as a plugin, mish is used in yolov4 | |prelu| mxnet's prelu activation with trainable gamma is implemented as a plugin, used in arcface | -|HardSwish| HardSwish activation is implemented as a plugin, used in yolov5 v3.0 | +|HardSwish| hard_swish = x * hard_sigmoid, used in yolov5 v3.0 | ## Speed Benchmark diff --git a/yolov5/CMakeLists.txt b/yolov5/CMakeLists.txt index 89f6d1c..d67668a 100644 --- a/yolov5/CMakeLists.txt +++ b/yolov5/CMakeLists.txt @@ -26,7 +26,7 @@ endif() set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -Wall -Ofast -Wfatal-errors -D_MWAITXINTRIN_H_INCLUDED") -cuda_add_library(myplugins SHARED ${PROJECT_SOURCE_DIR}/yololayer.cu ${PROJECT_SOURCE_DIR}/hardswish.cu) +cuda_add_library(myplugins SHARED ${PROJECT_SOURCE_DIR}/yololayer.cu) target_link_libraries(myplugins nvinfer cudart) find_package(OpenCV) diff --git a/yolov5/common.hpp b/yolov5/common.hpp index cb502a1..d31f7fc 100644 --- a/yolov5/common.hpp +++ b/yolov5/common.hpp @@ -9,7 +9,6 @@ #include #include "NvInfer.h" #include "yololayer.h" -#include "hardswish.h" #define CHECK(status) \ do\ @@ -201,13 +200,14 @@ ILayer* convBlock(INetworkDefinition *network, std::map& w conv1->setNbGroups(g); IScaleLayer* bn1 = addBatchNorm2d(network, weightMap, *conv1->getOutput(0), lname + ".bn", 1e-3); - auto creator = getPluginRegistry()->getPluginCreator("HardSwishLayer_TRT", "1"); - const PluginFieldCollection* pluginData = creator->getFieldNames(); - IPluginV2 *pluginObj = creator->createPlugin(("hardswish" + lname).c_str(), pluginData); - ITensor* inputTensors[] = {bn1->getOutput(0)}; - auto hs = network->addPluginV2(inputTensors, 1, *pluginObj); - - return hs; + // hard_swish = x * hard_sigmoid + auto hsig = network->addActivation(*bn1->getOutput(0), ActivationType::kHARD_SIGMOID); + assert(hsig); + hsig->setAlpha(1.0 / 6.0); + hsig->setBeta(0.5); + auto ew = network->addElementWise(*bn1->getOutput(0), *hsig->getOutput(0), ElementWiseOperation::kPROD); + assert(ew); + return ew; } ILayer* focus(INetworkDefinition *network, std::map& weightMap, ITensor& input, int inch, int outch, int ksize, std::string lname) { diff --git a/yolov5/hardswish.cu b/yolov5/hardswish.cu deleted file mode 100644 index 97719dd..0000000 --- a/yolov5/hardswish.cu +++ /dev/null @@ -1,182 +0,0 @@ -#include -#include "hardswish.h" -#include "utils.h" - -namespace nvinfer1 -{ - HardSwishPlugin::HardSwishPlugin() - { - } - - HardSwishPlugin::~HardSwishPlugin() - { - } - - // create the plugin at runtime from a byte stream - HardSwishPlugin::HardSwishPlugin(const void* data, size_t length) - { - const char *d = reinterpret_cast(data), *a = d; - Tn::read(d, mInputSize); - assert(d == a + length); - } - - void HardSwishPlugin::serialize(void* buffer) const - { - char* d = static_cast(buffer), *a = d; - Tn::write(d, mInputSize); - assert(d == a + getSerializationSize()); - } - - size_t HardSwishPlugin::getSerializationSize() const - { - return sizeof(mInputSize); - } - - int HardSwishPlugin::initialize() - { - return 0; - } - - Dims HardSwishPlugin::getOutputDimensions(int index, const Dims* inputs, int nbInputDims) - { - assert(nbInputDims == 1); - assert(index == 0); - return Dims3(inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]); - } - - // Set plugin namespace - void HardSwishPlugin::setPluginNamespace(const char* pluginNamespace) - { - mPluginNamespace = pluginNamespace; - } - - const char* HardSwishPlugin::getPluginNamespace() const - { - return mPluginNamespace; - } - - // Return the DataType of the plugin output at the requested index - DataType HardSwishPlugin::getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const - { - return DataType::kFLOAT; - } - - // Return true if output tensor is broadcast across a batch. - bool HardSwishPlugin::isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const - { - return false; - } - - // Return true if plugin can use input that is broadcast across batch without replication. - bool HardSwishPlugin::canBroadcastInputAcrossBatch(int inputIndex) const - { - return false; - } - - void HardSwishPlugin::configurePlugin(const PluginTensorDesc* in, int nbInput, const PluginTensorDesc* out, int nbOutput) - { - mInputSize = in[0].dims.d[0] * in[0].dims.d[1] * in[0].dims.d[2]; - } - - // Attach the plugin object to an execution context and grant the plugin the access to some context resource. - void HardSwishPlugin::attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) - { - } - - // Detach the plugin object from its execution context. - void HardSwishPlugin::detachFromContext() {} - - const char* HardSwishPlugin::getPluginType() const - { - return "HardSwishLayer_TRT"; - } - - const char* HardSwishPlugin::getPluginVersion() const - { - return "1"; - } - - void HardSwishPlugin::destroy() - { - delete this; - } - - // Clone the plugin - IPluginV2IOExt* HardSwishPlugin::clone() const - { - HardSwishPlugin *p = new HardSwishPlugin(); - p->setPluginNamespace(mPluginNamespace); - p->setInputSize(mInputSize); - return p; - } - - - __global__ void HardSwishKer(const float *in, float *out, int size) { - int idx = threadIdx.x + blockIdx.x * blockDim.x; - if (idx >= size) - return; - - if (in[idx] >= 3.0f) - out[idx] = in[idx]; - else if (in[idx] < -3.0f) - out[idx] = 0.0f; - else - out[idx] = in[idx] * (in[idx] + 3.0f) / 6.0f; - } - - void HardSwishPlugin::forwardGpu(const float *const * inputs, float* output, cudaStream_t stream, int batchSize) { - - int numElem = batchSize * mInputSize; - HardSwishKer<<<(numElem + mThreadCount - 1) / mThreadCount, mThreadCount>>> - (inputs[0], output, numElem); - } - - - int HardSwishPlugin::enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) - { - forwardGpu((const float *const *)inputs, (float*)outputs[0], stream, batchSize); - return 0; - } - - PluginFieldCollection HardSwishPluginCreator::mFC{}; - std::vector HardSwishPluginCreator::mPluginAttributes; - - HardSwishPluginCreator::HardSwishPluginCreator() - { - mPluginAttributes.clear(); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - const char* HardSwishPluginCreator::getPluginName() const - { - return "HardSwishLayer_TRT"; - } - - const char* HardSwishPluginCreator::getPluginVersion() const - { - return "1"; - } - - const PluginFieldCollection* HardSwishPluginCreator::getFieldNames() - { - return &mFC; - } - - IPluginV2IOExt* HardSwishPluginCreator::createPlugin(const char* name, const PluginFieldCollection* fc) - { - HardSwishPlugin* obj = new HardSwishPlugin(); - obj->setPluginNamespace(mNamespace.c_str()); - return obj; - } - - IPluginV2IOExt* HardSwishPluginCreator::deserializePlugin(const char* name, const void* serialData, size_t serialLength) - { - // This object will be deleted when the network is destroyed, which will - // call MishPlugin::destroy() - HardSwishPlugin* obj = new HardSwishPlugin(serialData, serialLength); - obj->setPluginNamespace(mNamespace.c_str()); - return obj; - } - -} diff --git a/yolov5/hardswish.h b/yolov5/hardswish.h deleted file mode 100644 index 700fa85..0000000 --- a/yolov5/hardswish.h +++ /dev/null @@ -1,87 +0,0 @@ -#ifndef _HARD_SWISH_LAYER_H -#define _HARD_SWISH_LAYER_H - -#include -#include -#include "NvInfer.h" - -namespace nvinfer1 -{ - class HardSwishPlugin: public IPluginV2IOExt - { - public: - explicit HardSwishPlugin(); - HardSwishPlugin(const void* data, size_t length); - ~HardSwishPlugin(); - - int getNbOutputs() const override - { - return 1; - } - - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override; - int initialize() override; - virtual void terminate() override {}; - virtual size_t getWorkspaceSize(int maxBatchSize) const override { return 0;} - virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override; - virtual size_t getSerializationSize() const override; - virtual void serialize(void* buffer) const override; - - bool supportsFormatCombination(int pos, const PluginTensorDesc* inOut, int nbInputs, int nbOutputs) const override { - return inOut[pos].format == TensorFormat::kLINEAR && inOut[pos].type == DataType::kFLOAT; - } - - const char* getPluginType() const override; - const char* getPluginVersion() const override; - void destroy() override; - IPluginV2IOExt* clone() const override; - void setPluginNamespace(const char* pluginNamespace) override; - const char* getPluginNamespace() const override; - DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const override; - bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const override; - bool canBroadcastInputAcrossBatch(int inputIndex) const override; - void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) override; - void configurePlugin(const PluginTensorDesc* in, int nbInput, const PluginTensorDesc* out, int nbOutput) override; - void detachFromContext() override; - - void setInputSize(int s) { - mInputSize = s; - } - - private: - void forwardGpu(const float *const * inputs,float * output, cudaStream_t stream,int batchSize = 1); - int mThreadCount = 256; - int mInputSize; - const char* mPluginNamespace; - }; - - class HardSwishPluginCreator : public IPluginCreator - { - public: - HardSwishPluginCreator(); - ~HardSwishPluginCreator() override = default; - const char* getPluginName() const override; - const char* getPluginVersion() const override; - const PluginFieldCollection* getFieldNames() override; - IPluginV2IOExt* createPlugin(const char* name, const PluginFieldCollection* fc) override; - IPluginV2IOExt* deserializePlugin(const char* name, const void* serialData, size_t serialLength) override; - - void setPluginNamespace(const char* libNamespace) override - { - mNamespace = libNamespace; - } - - const char* getPluginNamespace() const override - { - return mNamespace.c_str(); - } - - private: - std::string mNamespace; - static PluginFieldCollection mFC; - static std::vector mPluginAttributes; - }; - REGISTER_TENSORRT_PLUGIN(HardSwishPluginCreator); -}; - -#endif