haotian/duan8
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

yolov5 hswish plugin replaced by x*hsigmoid

Browse Source
This commit is contained in:
wang-xinyu 2020-09-10 19:43:30 +08:00
parent 2abc7b8375
commit 438b46e0dd
5 changed files with 10 additions and 279 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
README.md
View File

@ -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

2
yolov5/CMakeLists.txt
View File

@ -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)

16
yolov5/common.hpp
View File

@ -9,7 +9,6 @@
#include <dirent.h>
#include "NvInfer.h"
#include "yololayer.h"
#include "hardswish.h"
#define CHECK(status) \
do\
@ -201,13 +200,14 @@ ILayer* convBlock(INetworkDefinition *network, std::map<std::string, Weights>& 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<std::string, Weights>& weightMap, ITensor& input, int inch, int outch, int ksize, std::string lname) {

182
yolov5/hardswish.cu
View File

@ -1,182 +0,0 @@
#include <assert.h>
#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<const char *>(data), *a = d;
Tn::read(d, mInputSize);
assert(d == a + length);
}
void HardSwishPlugin::serialize(void* buffer) const
{
char* d = static_cast<char*>(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<PluginField> 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;
}
}

87
yolov5/hardswish.h
View File

@ -1,87 +0,0 @@
#ifndef _HARD_SWISH_LAYER_H
#define _HARD_SWISH_LAYER_H
#include <vector>
#include <string>
#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<PluginField> mPluginAttributes;
};
REGISTER_TENSORRT_PLUGIN(HardSwishPluginCreator);
};
#endif