mmr/tkDNN
Archived
1
0
Fork 0
Code Issues 46 Pull Requests 12 Releases 4 Activity

yoloRT load anchors

Browse Source
This commit is contained in:
Francesco Gatti
2019-02-18 21:39:14 +01:00
parent bdd8e0bc26
commit 87fe342ca2
18 changed files with 164 additions and 207 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/NetworkRT.cpp
+80 -100
View File
@@ -11,14 +11,6 @@
#include "NetworkRT.h"
using namespace nvinfer1;
#include "pluginsRT/ActivationLeakyRT.cpp"
#include "pluginsRT/ReorgRT.cpp"
#include "pluginsRT/RegionRT.cpp"
//#include "pluginsRT/RouteRT.cpp"
#include "pluginsRT/ShortcutRT.cpp"
#include "pluginsRT/YoloRT.cpp"
#include "pluginsRT/UpsampleRT.cpp"
#include "pluginsRT/Int8Calibrator.cpp"
// Logger for info/warning/errors
class Logger : public ILogger {
@@ -54,14 +46,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
builderRT->setMaxBatchSize(1);
builderRT->setMaxWorkspaceSize(1 << 30);
//change datatype based on system specs
if(builderRT->platformHasFastInt8()) {
BatchStream bstream({32,dim.c, dim.h, dim.w}, 32, 1);
Int8EntropyCalibrator calib(bstream, 0, false);
builderRT->setInt8Mode(true);
builderRT->setInt8Calibrator(&calib);
} else if(net->fp16 && builderRT->platformHasFastFp16()) {
if(net->fp16 && builderRT->platformHasFastFp16()) {
dtRT = DataType::kHALF;
builderRT->setHalf2Mode(true);
}
@@ -393,87 +378,6 @@ bool NetworkRT::serialize(const char *filename) {
return true;
}
class PluginFactory : IPluginFactory
{
public:
virtual IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength) {
const char * buf = reinterpret_cast<const char*>(serialData);
std::string name(layerName);
if(name.find("Activation") == 0) {
ActivationLeakyRT *a = new ActivationLeakyRT();
a->size = readBUF<int>(buf);
return a;
}
if(name.find("Region") == 0) {
RegionRT *r = new RegionRT(readBUF<int>(buf), //classes
readBUF<int>(buf), //coords
readBUF<int>(buf)); //num
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
return r;
}
if(name.find("Reorg") == 0) {
ReorgRT *r = new ReorgRT(readBUF<int>(buf)); //stride
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
return r;
}
if(name.find("Shortcut") == 0) {
ShortcutRT *r = new ShortcutRT();
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
return r;
}
if(name.find("Yolo") == 0) {
YoloRT *r = new YoloRT(readBUF<int>(buf), //classes
readBUF<int>(buf)); //num
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
for(int i=0; i<r->num; i++)
r->mask[i] = readBUF<dnnType>(buf);
for(int i=0; i<3*2*r->num; i++)
r->bias[i] = readBUF<dnnType>(buf);
std::cout<<"YOLO: "<<r->c<<" "<<r->h<<" "<<r->w<<"\n";
return r;
}
if(name.find("Upsample") == 0) {
UpsampleRT *r = new UpsampleRT(readBUF<int>(buf)); //stride
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
return r;
}
/*
if(name.find("Route") == 0) {
RouteRT *r = new RouteRT();
r->in = readBUF<int>(buf);
for(int i=0; i<RouteRT::MAX_INPUTS; i++)
r->c_in[i] = readBUF<int>(buf);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
return r;
}
*/
FatalError("Cant deserialize Plugin");
return NULL;
}
};
bool NetworkRT::deserialize(const char *filename) {
char *gieModelStream{nullptr};
@@ -488,13 +392,89 @@ bool NetworkRT::deserialize(const char *filename) {
file.close();
}
PluginFactory plfact;
runtimeRT = createInferRuntime(loggerRT);
engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) &plfact);
engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) pluginFactory);
//if (gieModelStream) delete [] gieModelStream;
return true;
}
IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) {
const char * buf = reinterpret_cast<const char*>(serialData);
std::string name(layerName);
if(name.find("Activation") == 0) {
ActivationLeakyRT *a = new ActivationLeakyRT();
a->size = readBUF<int>(buf);
return a;
}
if(name.find("Region") == 0) {
RegionRT *r = new RegionRT(readBUF<int>(buf), //classes
readBUF<int>(buf), //coords
readBUF<int>(buf)); //num
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
return r;
}
if(name.find("Reorg") == 0) {
ReorgRT *r = new ReorgRT(readBUF<int>(buf)); //stride
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
return r;
}
if(name.find("Shortcut") == 0) {
ShortcutRT *r = new ShortcutRT();
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
return r;
}
if(name.find("Yolo") == 0) {
YoloRT *r = new YoloRT(readBUF<int>(buf), //classes
readBUF<int>(buf)); //num
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
for(int i=0; i<r->num; i++)
r->mask[i] = readBUF<dnnType>(buf);
for(int i=0; i<3*2*r->num; i++)
r->bias[i] = readBUF<dnnType>(buf);
yolos[n_yolos++] = r;
return r;
}
if(name.find("Upsample") == 0) {
UpsampleRT *r = new UpsampleRT(readBUF<int>(buf)); //stride
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
return r;
}
/*
if(name.find("Route") == 0) {
RouteRT *r = new RouteRT();
r->in = readBUF<int>(buf);
for(int i=0; i<RouteRT::MAX_INPUTS; i++)
r->c_in[i] = readBUF<int>(buf);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
return r;
}
*/
FatalError("Cant deserialize Plugin");
return NULL;
}
}}
src/Yolo.cpp
+7 -44
View File
@@ -18,13 +18,12 @@ Yolo::Yolo(Network *net, int classes, int num, const char* fname_weights) :
this->num = num;
// load anchors
int seek = 0;
readBinaryFile(fname_weights, num, &mask_h, &mask_d, seek);
seek += num;
readBinaryFile(fname_weights, 3*num*2, &bias_h, &bias_d, seek);
printDeviceVector(num, mask_h, false);
printDeviceVector(3*num*2, bias_h, false);
if(fname_weights != nullptr) {
int seek = 0;
readBinaryFile(fname_weights, num, &mask_h, &mask_d, seek);
seek += num;
readBinaryFile(fname_weights, 3*num*2, &bias_h, &bias_d, seek);
}
// same
output_dim.n = input_dim.n;
@@ -33,10 +32,6 @@ Yolo::Yolo(Network *net, int classes, int num, const char* fname_weights) :
output_dim.w = input_dim.w;
output_dim.l = input_dim.l;
std::cout<<"YOLO INPUT: ";
input_dim.print();
std::cout<<"\n";
checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
predictions = nullptr;
}
@@ -62,35 +57,6 @@ Yolo::box get_yolo_box(float *x, float *biases, int n, int index, int i, int j,
return b;
}
void correct_yolo_boxes(Yolo::detection *dets, int n, int w, int h, int netw, int neth, int relative)
{
int i;
int new_w=0;
int new_h=0;
if (((float)netw/w) < ((float)neth/h)) {
new_w = netw;
new_h = (h * netw)/w;
} else {
new_h = neth;
new_w = (w * neth)/h;
}
for (i = 0; i < n; ++i){
Yolo::box b = dets[i].bbox;
b.x = (b.x - (netw - new_w)/2./netw) / ((float)new_w/netw);
b.y = (b.y - (neth - new_h)/2./neth) / ((float)new_h/neth);
b.w *= (float)netw/new_w;
b.h *= (float)neth/new_h;
if(!relative){
b.x *= w;
b.w *= w;
b.y *= h;
b.h *= h;
}
dets[i].bbox = b;
}
}
dnnType* Yolo::infer(dataDim_t &dim, dnnType* srcData) {
checkCuda( cudaMemcpy(dstData, srcData, dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice));
@@ -109,14 +75,12 @@ dnnType* Yolo::infer(dataDim_t &dim, dnnType* srcData) {
return dstData;
}
int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int w, int h, int netw, int neth, float thresh) {
int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh) {
if(predictions == nullptr)
predictions = new dnnType[output_dim.tot()];
checkCuda( cudaMemcpy(predictions, dstData, output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost));
int relative = 0;
int lw = output_dim.w;
int lh = output_dim.h;
@@ -150,7 +114,6 @@ int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int w, int h, int
}
}
correct_yolo_boxes(dets + ndets, count, w, h, netw, neth, relative);
ndets = count;
return count;
}
src/Yolo3Detection.cpp
+28 -19
View File
@@ -2,10 +2,36 @@
namespace tk { namespace dnn {
bool Yolo3Detection::init(std::string tensor_folder) {
bool Yolo3Detection::init(std::string tensor_path) {
//const char *tensor_path = "../data/yolo3/yolo3_berkeley.rt";
//convert network to tensorRT
std::cout<<(tensor_path).c_str()<<"\n";
netRT = new tk::dnn::NetworkRT(NULL, (tensor_path).c_str() );
if(netRT->pluginFactory->n_yolos != 3) {
FatalError("this is not yolo3");
}
for(int i=0; i<netRT->pluginFactory->n_yolos; i++) {
YoloRT *yRT = netRT->pluginFactory->yolos[i];
classes = yRT->classes;
num = yRT->num;
// make a yolo layer for interpret predictions
yolo[i] = new tk::dnn::Yolo(nullptr, classes, num, nullptr); // yolo without input and bias
memcpy(yolo[i]->mask_h, yRT->mask, sizeof(dnnType)*num);
memcpy(yolo[i]->bias_h, yRT->bias, sizeof(dnnType)*num*3*2);
yolo[i]->input_dim = yolo[i]->output_dim = tk::dnn::dataDim_t(1, yRT->c, yRT->h, yRT->w);
}
dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes);
checkCuda(cudaMallocHost(&input, sizeof(dnnType)*netRT->input_dim.tot()));
checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot()));
// class colors precompute
for(int c=0; c<classes; c++) {
int cc = c+1;
@@ -19,23 +45,6 @@ bool Yolo3Detection::init(std::string tensor_folder) {
//std::cout<<r<<" "<<g<<" "<<b<<"\n";
colors[c] = cv::Scalar(int(255.0*b), int(255.0*g), int(255.0*r));
}
//convert network to tensorRT
std::cout<<(tensor_folder + ".rt").c_str()<<"\n";
netRT = new tk::dnn::NetworkRT(NULL, (tensor_folder + ".rt").c_str() );
yolo[0] = new tk::dnn::Yolo(nullptr, classes, num, (tensor_folder + "_0.bin").c_str() ); // yolo without input and bias
yolo[0]->input_dim = yolo[0]->output_dim = tk::dnn::dataDim_t(1, 45, 10, 17);
yolo[1] = new tk::dnn::Yolo(nullptr, classes, num, (tensor_folder + "_1.bin").c_str() ); // yolo without input and bias
yolo[1]->input_dim = yolo[1]->output_dim = tk::dnn::dataDim_t(1, 45, 20, 34);
yolo[2] = new tk::dnn::Yolo(nullptr, classes, num, (tensor_folder + "_2.bin").c_str() ); // yolo without input and bias
yolo[2]->input_dim = yolo[2]->output_dim = tk::dnn::dataDim_t(1, 45, 40, 68);
dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes);
checkCuda(cudaMallocHost(&input, sizeof(dnnType)*netRT->input_dim.tot()));
checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot()));
return true;
}
@@ -82,7 +91,7 @@ void Yolo3Detection::update(cv::Mat &imageORIG) {
for(int i=0; i<3; i++) {
rt_out[i] = (dnnType*)netRT->buffersRT[i+1];
yolo[i]->dstData = rt_out[i];
yolo[i]->computeDetections(dets, ndets, netRT->input_dim.w, netRT->input_dim.h, netRT->input_dim.w, netRT->input_dim.h, thresh);
yolo[i]->computeDetections(dets, ndets, netRT->input_dim.w, netRT->input_dim.h, thresh);
}
tk::dnn::Yolo::mergeDetections(dets, ndets, classes);
TIMER_STOP
src/pluginsRT/ActivationLeakyRT.cpp
-60
View File
@@ -1,60 +0,0 @@
#include<cassert>
#include "kernels.h"
class ActivationLeakyRT : public IPlugin {
public:
ActivationLeakyRT() {
}
~ActivationLeakyRT(){
}
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<outputDims[0].nbDims; i++)
size *= outputDims[0].d[i];
}
int initialize() override {
return 0;
}
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 {
activationLEAKYForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), size, stream);
return 0;
}
virtual size_t getSerializationSize() override {
return 1*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, size);
}
int size;
};
src/pluginsRT/Int8Calibrator.cpp
-168
View File
@@ -1,168 +0,0 @@
#include <vector>
#include <assert.h>
#include <algorithm>
#include <iterator>
#include "NvInfer.h"
class BatchStream
{
public:
BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches)
{
mBatchSize = batchSize;
mMaxBatches = maxBatches;
mDims = nvinfer1::DimsNCHW{ dim.n, dim.c, dim.h, dim.w };
mImageSize = mDims.c()*mDims.h()*mDims.w();
mBatch.resize(mBatchSize*mImageSize, 0);
mLabels.resize(mBatchSize, 0);
mFileBatch.resize(mDims.n()*mImageSize, 0);
mFileLabels.resize(mDims.n(), 0);
reset(0);
}
void reset(int firstBatch)
{
mBatchCount = 0;
mFileCount = 0;
mFileBatchPos = mDims.n();
skip(firstBatch);
}
bool next()
{
std::cout<<"Next batch: "<<mBatchCount<<" of "<<mMaxBatches<<"\n";
if (mBatchCount == mMaxBatches)
return false;
for (int csize = 1, batchPos = 0; batchPos < mBatchSize; batchPos += csize, mFileBatchPos += csize)
{
assert(mFileBatchPos > 0 && mFileBatchPos <= mDims.n());
if (mFileBatchPos == mDims.n() && !update())
return false;
// copy the smaller of: elements left to fulfill the request, or elements left in the file buffer.
csize = std::min(mBatchSize - batchPos, mDims.n() - mFileBatchPos);
std::copy_n(getFileBatch() + mFileBatchPos * mImageSize, csize * mImageSize, getBatch() + batchPos * mImageSize);
std::copy_n(getFileLabels() + mFileBatchPos, csize, getLabels() + batchPos);
}
mBatchCount++;
return true;
}
void skip(int skipCount)
{
if (mBatchSize >= mDims.n() && mBatchSize%mDims.n() == 0 && mFileBatchPos == mDims.n())
{
mFileCount += skipCount * mBatchSize / mDims.n();
std::cout<<mFileCount<<"\n";
return;
}
int x = mBatchCount;
for (int i = 0; i < skipCount; i++)
next();
mBatchCount = x;
}
float *getBatch() { return &mBatch[0]; }
float *getLabels() { return &mLabels[0]; }
int getBatchesRead() const { return mBatchCount; }
int getBatchSize() const { return mBatchSize; }
nvinfer1::DimsNCHW getDims() const { return mDims; }
private:
float* getFileBatch() { return &mFileBatch[0]; }
float* getFileLabels() { return &mFileLabels[0]; }
bool update()
{
std::string inputFileName = std::string("calibBatches/batch") + std::to_string(mFileCount++);
FILE * file = fopen(inputFileName.c_str(), "rb");
if (!file) {
FatalError("cant open batch calib file: " + inputFileName);
return false;
}
size_t readInputCount = fread(getFileBatch(), sizeof(float), mDims.n()*mImageSize, file);
size_t readLabelCount = fread(getFileLabels(), sizeof(float), mDims.n(), file);;
assert(readInputCount == size_t(mDims.n()*mImageSize) && readLabelCount == size_t(mDims.n()));
fclose(file);
mFileBatchPos = 0;
return true;
}
int mBatchSize{ 0 };
int mMaxBatches{ 0 };
int mBatchCount{ 0 };
int mFileCount{ 0 }, mFileBatchPos{ 0 };
int mImageSize{ 0 };
nvinfer1::DimsNCHW mDims;
std::vector<float> mBatch;
std::vector<float> mLabels;
std::vector<float> mFileBatch;
std::vector<float> mFileLabels;
};
class Int8EntropyCalibrator : public IInt8EntropyCalibrator
{
public:
Int8EntropyCalibrator(BatchStream& stream, int firstBatch, bool readCache = true)
: mStream(stream), mReadCache(readCache)
{
DimsNCHW dims = mStream.getDims();
mInputCount = mStream.getBatchSize() * dims.c() * dims.h() * dims.w();
checkCuda(cudaMalloc(&mDeviceInput, mInputCount * sizeof(float)));
mStream.reset(firstBatch);
}
virtual ~Int8EntropyCalibrator()
{
checkCuda(cudaFree(mDeviceInput));
}
int getBatchSize() const override { return mStream.getBatchSize(); }
bool getBatch(void* bindings[], const char* names[], int nbBindings) override
{
std::cout<<"CALIB request batch\n";
if (!mStream.next())
return false;
checkCuda(cudaMemcpy(mDeviceInput, mStream.getBatch(), mInputCount * sizeof(float), cudaMemcpyHostToDevice));
bindings[0] = mDeviceInput;
return true;
}
const void* readCalibrationCache(size_t& length) override
{
mCalibrationCache.clear();
std::ifstream input("table.calib", std::ios::binary);
input >> std::noskipws;
FatalError("rewrite different");
//if (mReadCache && input.good())
// std::copy(std::istream_iterator<char>(input), std::istream_iterator<char>(), std::back_inserter(mCalibrationCache));
length = mCalibrationCache.size();
return length ? &mCalibrationCache[0] : nullptr;
}
void writeCalibrationCache(const void* cache, size_t length) override
{
std::ofstream output("table.calib", std::ios::binary);
output.write(reinterpret_cast<const char*>(cache), length);
}
private:
BatchStream mStream;
bool mReadCache{ true };
size_t mInputCount;
void* mDeviceInput{ nullptr };
std::vector<char> mCalibrationCache;
};
src/pluginsRT/RegionRT.cpp
-94
View File
@@ -1,94 +0,0 @@
#include<cassert>
#include "kernels.h"
class RegionRT : public IPlugin {
public:
RegionRT(int classes, int coords, int num) {
this->classes = classes;
this->coords = coords;
this->num = num;
}
~RegionRT(){
}
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 {
c = inputDims[0].d[0];
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
int initialize() override {
return 0;
}
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 {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
for (int b = 0; b < batchSize; ++b){
for(int n = 0; n < num; ++n){
int index = entry_index(b, n*w*h, 0, batchSize);
activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream);
index = entry_index(b, n*w*h, coords, batchSize);
activationLOGISTICForward(srcData + index, dstData + index, w*h, stream);
}
}
//softmax start
int index = entry_index(0, 0, coords + 1, batchSize);
softmaxForward( srcData + index, classes, batchSize*num,
(batchSize*c*h*w)/num,
w*h, 1, w*h, 1, dstData + index, stream);
return 0;
}
virtual size_t getSerializationSize() override {
return 6*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, classes);
tk::dnn::writeBUF(buf, coords);
tk::dnn::writeBUF(buf, num);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
}
int c, h, w;
int classes, coords, num;
int entry_index(int batch, int location, int entry, int batchSize) {
int n = location / (w*h);
int loc = location % (w*h);
return batch*c*h*w*batchSize + n*w*h*(coords+classes+1) + entry*w*h + loc;
}
};
src/pluginsRT/ReorgRT.cpp
-63
View File
@@ -1,63 +0,0 @@
#include<cassert>
#include "kernels.h"
class ReorgRT : public IPlugin {
public:
ReorgRT(int stride) {
this->stride = stride;
}
~ReorgRT(){
}
int getNbOutputs() const override {
return 1;
}
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
return DimsCHW{inputs[0].d[0]*stride*stride, inputs[0].d[1]/stride, inputs[0].d[2]/stride};
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
c = inputDims[0].d[0];
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
int initialize() override {
return 0;
}
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 {
reorgForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]),
batchSize, c, h, w, stride, stream);
return 0;
}
virtual size_t getSerializationSize() override {
return 4*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, stride);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
}
int c, h, w, stride;
};
src/pluginsRT/RouteRT.cpp
-82
View File
@@ -1,82 +0,0 @@
#include<cassert>
#include "kernels.h"
class RouteRT : public IPlugin {
public:
RouteRT() {
}
~RouteRT(){
}
int getNbOutputs() const override {
return 1;
}
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
int out_c = 0;
for(int i=0; i<nbInputDims; i++) out_c += inputs[i].d[0];
return DimsCHW{out_c, inputs[0].d[1], inputs[0].d[2]};
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
in = nbInputs;
c = 0;
for(int i=0; i<nbInputs; i++) {
c_in[i] = inputDims[i].d[0];
c += inputDims[i].d[0];
}
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
int initialize() override {
return 0;
}
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 {
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
int offset = 0;
for(int i=0; i<in; i++) {
dnnType *input = (dnnType*)reinterpret_cast<const dnnType*>(inputs[i]);
int in_dim = c_in[i]*h*w;
checkCuda( cudaMemcpyAsync(dstData + offset, input, in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) );
offset += in_dim;
}
return 0;
}
virtual size_t getSerializationSize() override {
return (4+MAX_INPUTS)*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, in);
for(int i=0; i<MAX_INPUTS; i++)
tk::dnn::writeBUF(buf, c_in[i]);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
}
static const int MAX_INPUTS = 4;
int in;
int c_in[MAX_INPUTS];
int c, h, w;
};
src/pluginsRT/ShortcutRT.cpp
-65
View File
@@ -1,65 +0,0 @@
#include<cassert>
#include "kernels.h"
class ShortcutRT : public IPlugin {
public:
ShortcutRT() {
}
~ShortcutRT(){
}
int getNbOutputs() const override {
return 1;
}
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
return DimsCHW{inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]};
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
c = inputDims[0].d[0];
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
int initialize() override {
return 0;
}
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 {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *srcDataBack = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, c, h, w, 1, stream);
return 0;
}
virtual size_t getSerializationSize() override {
return 3*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
}
int c, h, w;
};
src/pluginsRT/UpsampleRT.cpp
-65
View File
@@ -1,65 +0,0 @@
#include<cassert>
#include "kernels.h"
class UpsampleRT : public IPlugin {
public:
UpsampleRT(int stride) {
this->stride = stride;
}
~UpsampleRT(){
}
int getNbOutputs() const override {
return 1;
}
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
return DimsCHW(inputs[0].d[0], inputs[0].d[1]*stride, inputs[0].d[2]*stride);
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
c = inputDims[0].d[0];
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
int initialize() override {
return 0;
}
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 {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
fill(dstData, batchSize*c*h*w*stride*stride, 0.0, stream);
upsampleForward(srcData, dstData, batchSize, c, h, w, stride, 1, 1, stream);
return 0;
}
virtual size_t getSerializationSize() override {
return 4*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, stride);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
}
int c, h, w, stride;
};
src/pluginsRT/YoloRT.cpp
-103
View File
@@ -1,103 +0,0 @@
#include<cassert>
#include "kernels.h"
class YoloRT : public IPlugin {
public:
YoloRT(int classes, int num, tk::dnn::Yolo *yolo = nullptr) {
this->classes = classes;
this->num = num;
mask = new dnnType[num];
bias = new dnnType[num*3*2];
if(yolo != nullptr) {
memcpy(mask, yolo->mask_h, sizeof(dnnType)*num);
memcpy(bias, yolo->bias_h, sizeof(dnnType)*num*3*2);
}
}
~YoloRT(){
}
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 {
c = inputDims[0].d[0];
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
int initialize() override {
return 0;
}
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 {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
for (int b = 0; b < batchSize; ++b){
for(int n = 0; n < num; ++n){
int index = entry_index(b, n*w*h, 0, batchSize);
activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream);
index = entry_index(b, n*w*h, 4, batchSize);
activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream);
}
}
//std::cout<<"YOLO END\n";
return 0;
}
virtual size_t getSerializationSize() override {
return 5*sizeof(int) + num*sizeof(dnnType) + num*3*2*sizeof(dnnType);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, classes);
tk::dnn::writeBUF(buf, num);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
for(int i=0; i<num; i++)
tk::dnn::writeBUF(buf, mask[i]);
for(int i=0; i<3*2*num; i++)
tk::dnn::writeBUF(buf, bias[i]);
}
int c, h, w;
int classes, num;
dnnType *mask;
dnnType *bias;
int entry_index(int batch, int location, int entry, int batchSize) {
int n = location / (w*h);
int loc = location % (w*h);
return batch*c*h*w*batchSize + n*w*h*(4+classes+1) + entry*w*h + loc;
}
};