Add swish #249
@@ -19,6 +19,7 @@ enum layerType_t {
|
|||||||
LAYER_ACTIVATION_CRELU,
|
LAYER_ACTIVATION_CRELU,
|
||||||
LAYER_ACTIVATION_LEAKY,
|
LAYER_ACTIVATION_LEAKY,
|
||||||
LAYER_ACTIVATION_MISH,
|
LAYER_ACTIVATION_MISH,
|
||||||
|
LAYER_ACTIVATION_SWISH,
|
||||||
LAYER_ACTIVATION_LOGISTIC,
|
LAYER_ACTIVATION_LOGISTIC,
|
||||||
LAYER_FLATTEN,
|
LAYER_FLATTEN,
|
||||||
LAYER_RESHAPE,
|
LAYER_RESHAPE,
|
||||||
@@ -75,6 +76,7 @@ public:
|
|||||||
case LAYER_ACTIVATION_CRELU: return "ActivationCReLU";
|
case LAYER_ACTIVATION_CRELU: return "ActivationCReLU";
|
||||||
case LAYER_ACTIVATION_LEAKY: return "ActivationLeaky";
|
case LAYER_ACTIVATION_LEAKY: return "ActivationLeaky";
|
||||||
case LAYER_ACTIVATION_MISH: return "ActivationMish";
|
case LAYER_ACTIVATION_MISH: return "ActivationMish";
|
||||||
|
case LAYER_ACTIVATION_SWISH: return "ActivationSwish";
|
||||||
case LAYER_ACTIVATION_LOGISTIC: return "ActivationLogistic";
|
case LAYER_ACTIVATION_LOGISTIC: return "ActivationLogistic";
|
||||||
case LAYER_FLATTEN: return "Flatten";
|
case LAYER_FLATTEN: return "Flatten";
|
||||||
case LAYER_RESHAPE: return "Reshape";
|
case LAYER_RESHAPE: return "Reshape";
|
||||||
@@ -223,7 +225,8 @@ typedef enum {
|
|||||||
ACTIVATION_ELU = 100,
|
ACTIVATION_ELU = 100,
|
||||||
ACTIVATION_LEAKY = 101,
|
ACTIVATION_LEAKY = 101,
|
||||||
ACTIVATION_MISH = 102,
|
ACTIVATION_MISH = 102,
|
||||||
ACTIVATION_LOGISTIC = 103
|
ACTIVATION_LOGISTIC = 103,
|
||||||
|
ACTIVATION_SWISH = 104
|
||||||
} tkdnnActivationMode_t;
|
} tkdnnActivationMode_t;
|
||||||
|
|
||||||
/**
|
/**
|
||||||
@@ -245,6 +248,8 @@ public:
|
|||||||
return LAYER_ACTIVATION_LEAKY;
|
return LAYER_ACTIVATION_LEAKY;
|
||||||
else if (act_mode == ACTIVATION_MISH)
|
else if (act_mode == ACTIVATION_MISH)
|
||||||
return LAYER_ACTIVATION_MISH;
|
return LAYER_ACTIVATION_MISH;
|
||||||
|
else if (act_mode == ACTIVATION_SWISH)
|
||||||
|
return LAYER_ACTIVATION_SWISH;
|
||||||
else if (act_mode == ACTIVATION_LOGISTIC)
|
else if (act_mode == ACTIVATION_LOGISTIC)
|
||||||
return LAYER_ACTIVATION_LOGISTIC;
|
return LAYER_ACTIVATION_LOGISTIC;
|
||||||
else
|
else
|
||||||
|
|||||||
+63
-62
@@ -7,10 +7,11 @@
|
|||||||
#include "Layer.h"
|
#include "Layer.h"
|
||||||
#include "NvInfer.h"
|
#include "NvInfer.h"
|
||||||
#include <memory>
|
#include <memory>
|
||||||
#include <tkDNN/kernels.h>
|
#include <kernels.h>
|
||||||
#include <pluginsRT/ActivationLeakyRT.h>
|
#include <pluginsRT/ActivationLeakyRT.h>
|
||||||
#include <pluginsRT/ActivationLogisticRT.h>
|
#include <pluginsRT/ActivationLogisticRT.h>
|
||||||
#include <pluginsRT/ActivationMishRT.h>
|
#include <pluginsRT/ActivationMishRT.h>
|
||||||
|
#include <pluginsRT/ActivationSwishRT.h>
|
||||||
#include <pluginsRT/ActivationReLUCeilingRT.h>
|
#include <pluginsRT/ActivationReLUCeilingRT.h>
|
||||||
#include <pluginsRT/DeformableConvRT.h>
|
#include <pluginsRT/DeformableConvRT.h>
|
||||||
#include <pluginsRT/FlattenConcatRT.h>
|
#include <pluginsRT/FlattenConcatRT.h>
|
||||||
@@ -30,85 +31,85 @@
|
|||||||
|
|
||||||
namespace tk { namespace dnn {
|
namespace tk { namespace dnn {
|
||||||
|
|
||||||
class NetworkRT {
|
class NetworkRT {
|
||||||
|
|
||||||
public:
|
public:
|
||||||
nvinfer1::DataType dtRT;
|
nvinfer1::DataType dtRT;
|
||||||
nvinfer1::IBuilder *builderRT;
|
nvinfer1::IBuilder *builderRT;
|
||||||
nvinfer1::IRuntime *runtimeRT;
|
nvinfer1::IRuntime *runtimeRT;
|
||||||
nvinfer1::INetworkDefinition *networkRT;
|
nvinfer1::INetworkDefinition *networkRT;
|
||||||
#if NV_TENSORRT_MAJOR >= 6
|
#if NV_TENSORRT_MAJOR >= 6
|
||||||
nvinfer1::IBuilderConfig *configRT;
|
nvinfer1::IBuilderConfig *configRT;
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
nvinfer1::ICudaEngine *engineRT;
|
nvinfer1::ICudaEngine *engineRT;
|
||||||
nvinfer1::IExecutionContext *contextRT;
|
nvinfer1::IExecutionContext *contextRT;
|
||||||
|
|
||||||
const static int MAX_BUFFERS_RT = 10;
|
const static int MAX_BUFFERS_RT = 10;
|
||||||
void* buffersRT[MAX_BUFFERS_RT];
|
void* buffersRT[MAX_BUFFERS_RT];
|
||||||
dataDim_t buffersDIM[MAX_BUFFERS_RT];
|
dataDim_t buffersDIM[MAX_BUFFERS_RT];
|
||||||
int buf_input_idx, buf_output_idx;
|
int buf_input_idx, buf_output_idx;
|
||||||
bool builderActive = false;
|
bool builderActive = false;
|
||||||
dataDim_t input_dim, output_dim;
|
dataDim_t input_dim, output_dim;
|
||||||
dnnType *output;
|
dnnType *output;
|
||||||
cudaStream_t stream;
|
cudaStream_t stream;
|
||||||
|
|
||||||
std::vector<nvinfer1::YoloRT*> yolo_plugins; // yolo layers in network
|
std::vector<nvinfer1::YoloRT*> yolo_plugins; // yolo layers in network
|
||||||
|
|
||||||
NetworkRT(Network *net, const char *name);
|
NetworkRT(Network *net, const char *name);
|
||||||
virtual ~NetworkRT();
|
virtual ~NetworkRT();
|
||||||
|
|
||||||
int getMaxBatchSize() {
|
int getMaxBatchSize() {
|
||||||
if(engineRT != nullptr)
|
if(engineRT != nullptr)
|
||||||
return engineRT->getMaxBatchSize();
|
return engineRT->getMaxBatchSize();
|
||||||
else
|
else
|
||||||
return 0;
|
return 0;
|
||||||
}
|
}
|
||||||
|
|
||||||
int getBuffersN() {
|
int getBuffersN() {
|
||||||
if(engineRT != nullptr)
|
if(engineRT != nullptr)
|
||||||
return engineRT->getNbBindings();
|
return engineRT->getNbBindings();
|
||||||
else
|
else
|
||||||
return 0;
|
return 0;
|
||||||
}
|
}
|
||||||
|
|
||||||
/**
|
/**
|
||||||
Do inference
|
Do inference
|
||||||
*/
|
*/
|
||||||
dnnType* infer(dataDim_t &dim, dnnType* data);
|
dnnType* infer(dataDim_t &dim, dnnType* data);
|
||||||
void enqueue(int batchSize = 1);
|
void enqueue(int batchSize = 1);
|
||||||
|
|
||||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Layer *l);
|
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, Conv2d *l);
|
||||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Activation *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, Dense *l);
|
||||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Pooling *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, Softmax *l);
|
||||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Route *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, Flatten *l);
|
||||||
nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Reshape *l);
|
nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Reshape *l);
|
||||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Resize *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, Reorg *l);
|
||||||
nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Region *l);
|
nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Region *l);
|
||||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Shortcut *l);
|
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Shortcut *l);
|
||||||
nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Yolo *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, Upsample *l);
|
||||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, DeformConv2d *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,Padding *l);
|
||||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor* input,MulAdd *l);
|
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor* input,MulAdd *l);
|
||||||
|
|
||||||
#if NV_TENSORRT_MAJOR > 5 && NV_TENSORRT_MAJOR < 8
|
#if NV_TENSORRT_MAJOR > 5 && NV_TENSORRT_MAJOR < 8
|
||||||
bool serialize(const char *filename);
|
bool serialize(const char *filename);
|
||||||
#else
|
#else
|
||||||
bool serialize(const char *filename,nvinfer1::IHostMemory *ptr);
|
bool serialize(const char *filename,nvinfer1::IHostMemory *ptr);
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
bool deserialize(const char *filename);
|
bool deserialize(const char *filename);
|
||||||
void destroy();
|
void destroy();
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
};
|
};
|
||||||
|
|
||||||
}}
|
}}
|
||||||
#endif //NETWORKRT_H
|
#endif //NETWORKRT_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 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 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 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));
|
void fill(dnnType *data, int size, dnnType val, cudaStream_t stream = cudaStream_t(0));
|
||||||
|
|
||||||
|
|||||||
@@ -0,0 +1,82 @@
|
|||||||
|
#include<cassert>
|
||||||
|
#include "../kernels.h"
|
||||||
|
#include <NvInfer.h>
|
||||||
|
#include <vector>
|
||||||
|
|
||||||
|
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<PluginField> mPluginAttributes;
|
||||||
|
std::string mPluginNamespace;
|
||||||
|
};
|
||||||
|
|
||||||
|
REGISTER_TENSORRT_PLUGIN(ActivationSwishRTPluginCreator);
|
||||||
|
};
|
||||||
@@ -52,6 +52,10 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) {
|
|||||||
else if(act_mode == ACTIVATION_MISH) {
|
else if(act_mode == ACTIVATION_MISH) {
|
||||||
activationMishForward(srcData, dstData, dim.tot());
|
activationMishForward(srcData, dstData, dim.tot());
|
||||||
|
|
||||||
|
}
|
||||||
|
else if(act_mode == ACTIVATION_SWISH) {
|
||||||
|
activationSwishForward(srcData, dstData, dim.tot());
|
||||||
|
|
||||||
}
|
}
|
||||||
else if(act_mode == ACTIVATION_LOGISTIC) {
|
else if(act_mode == ACTIVATION_LOGISTIC) {
|
||||||
activationLOGISTICForward(srcData, dstData, dim.tot());
|
activationLOGISTICForward(srcData, dstData, dim.tot());
|
||||||
|
|||||||
@@ -197,6 +197,7 @@ namespace tk { namespace dnn {
|
|||||||
if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU);
|
if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU);
|
||||||
else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY;
|
else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY;
|
||||||
else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH;
|
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 if(f.activation == "logistic") act = tk::dnn::ACTIVATION_LOGISTIC;
|
||||||
else { FatalError("activation not supported: " + f.activation); }
|
else { FatalError("activation not supported: " + f.activation); }
|
||||||
netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act);
|
netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act);
|
||||||
|
|||||||
+15
-12
@@ -254,7 +254,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
|
|||||||
return convert_layer(input, (Conv2d*) l);
|
return convert_layer(input, (Conv2d*) l);
|
||||||
if(type == LAYER_POOLING)
|
if(type == LAYER_POOLING)
|
||||||
return convert_layer(input, (Pooling*) l);
|
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);
|
return convert_layer(input, (Activation*) l);
|
||||||
if(type == LAYER_SOFTMAX)
|
if(type == LAYER_SOFTMAX)
|
||||||
return convert_layer(input, (Softmax*) l);
|
return convert_layer(input, (Softmax*) l);
|
||||||
@@ -646,8 +646,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
|
|||||||
checkNULL(lRT);
|
checkNULL(lRT);
|
||||||
return lRT;
|
return lRT;
|
||||||
}
|
}
|
||||||
else if(l->act_mode == CUDNN_ACTIVATION_ELU || l->act_mode == ACTIVATION_ELU){
|
else if(l->act_mode == CUDNN_ACTIVATION_ELU || l->act_mode == ACTIVATION_ELU) {
|
||||||
IActivationLayer *lRT = networkRT->addActivation(*input,ActivationType::kELU);
|
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);
|
checkNULL(lRT);
|
||||||
return lRT;
|
return lRT;
|
||||||
}
|
}
|
||||||
@@ -1025,18 +1029,17 @@ bool NetworkRT::deserialize(const char *filename) {
|
|||||||
}
|
}
|
||||||
|
|
||||||
#if NV_TENSORRT_MAJOR > 7
|
#if NV_TENSORRT_MAJOR > 7
|
||||||
void NetworkRT::destroy() {
|
void NetworkRT::destroy() {
|
||||||
delete contextRT;
|
delete contextRT;
|
||||||
if(builderActive) {
|
if(builderActive) {
|
||||||
delete engineRT;
|
delete engineRT;
|
||||||
delete builderRT;
|
delete builderRT;
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
#elif NV_TENSORRT_MAJOR <=7
|
#elif NV_TENSORRT_MAJOR <=7
|
||||||
void NetworkRT::destroy() {
|
void NetworkRT::destroy() {
|
||||||
|
|
||||||
}
|
}
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
|
|
||||||
}}
|
}}
|
||||||
@@ -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<size) {
|
||||||
|
output[i] = input[i] * 1.0f/(1.0f + exp(-input[i]));;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
/**
|
||||||
|
Swish activation function
|
||||||
|
*/
|
||||||
|
void activationSwishForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream)
|
||||||
|
{
|
||||||
|
int blocks = (size+255)/256;
|
||||||
|
int threads = 256;
|
||||||
|
|
||||||
|
activation_swish<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
|
||||||
|
}
|
||||||
@@ -0,0 +1,133 @@
|
|||||||
|
//
|
||||||
|
// Created by Adam on 4/11/2022
|
||||||
|
//
|
||||||
|
#include <tkDNN/pluginsRT/ActivationSwishRT.h>
|
||||||
|
using namespace nvinfer1;
|
||||||
|
std::vector<PluginField> ActivationSwishRTPluginCreator::mPluginAttributes;
|
||||||
|
PluginFieldCollection ActivationSwishRTPluginCreator::mFC{};
|
||||||
|
|
||||||
|
ActivationSwishRT::ActivationSwishRT() {
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
ActivationSwishRT::~ActivationSwishRT() {
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
ActivationSwishRT::ActivationSwishRT(const void *data, size_t length) {
|
||||||
|
const char *buf = reinterpret_cast<const char *>(data), *bufCheck = buf;
|
||||||
|
size = readBUF<int>(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<const dnnType *>(inputs[0]),
|
||||||
|
reinterpret_cast<dnnType *>(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<const dnnType *>(inputs[0]),
|
||||||
|
reinterpret_cast<dnnType *>(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<char *>(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;
|
||||||
|
}
|
||||||
|
|
||||||
File diff suppressed because it is too large
Load Diff
@@ -0,0 +1,40 @@
|
|||||||
|
//
|
||||||
|
// Created by Adam T. Cuellar on 9/15/21.
|
||||||
|
//
|
||||||
|
|
||||||
|
#include<iostream>
|
||||||
|
#include<vector>
|
||||||
|
#include "tkdnn.h"
|
||||||
|
#include "test.h"
|
||||||
|
#include "DarknetParser.h"
|
||||||
|
|
||||||
|
int main() {
|
||||||
|
std::string bin_path = "yolov4-csp-swish";
|
||||||
|
std::vector<std::string> input_bins = {
|
||||||
|
bin_path + "/layers/input.bin"
|
||||||
|
};
|
||||||
|
std::vector<std::string> 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;
|
||||||
|
}
|
||||||
Reference in New Issue
Block a user