include dir fix, cmake dir

include/tkDNN/pluginsRT/ActivationLeakyRT.h

@@ -0,0 +1,60 @@
+#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;
+};

include/tkDNN/pluginsRT/Int8Calibrator.h

@@ -0,0 +1,168 @@
+#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;
+};

include/tkDNN/pluginsRT/RegionRT.h

@@ -0,0 +1,94 @@
+#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;
+	}
+
+};

include/tkDNN/pluginsRT/ReorgRT.h

@@ -0,0 +1,63 @@
+#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;
+};

include/tkDNN/pluginsRT/RouteRT.h

@@ -0,0 +1,82 @@
+#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;
+};

include/tkDNN/pluginsRT/ShortcutRT.h

@@ -0,0 +1,65 @@
+#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;
+};

include/tkDNN/pluginsRT/UpsampleRT.h

@@ -0,0 +1,65 @@
+#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;
+};

include/tkDNN/pluginsRT/YoloRT.h

@@ -0,0 +1,116 @@
+#include<cassert>
+#include "../kernels.h"
+
+#define YOLORT_CLASSNAME_W 256
+
+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);
+			classesNames = yolo->classesNames;
+        }
+	}
+
+	~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) + YOLORT_CLASSNAME_W*classes*sizeof(char);
+	}
+
+	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]);
+
+		// save classes names
+		for(int i=0; i<classes; i++) {
+			char tmp[YOLORT_CLASSNAME_W];
+			strcpy(tmp, classesNames[i].c_str());
+			for(int j=0; j<YOLORT_CLASSNAME_W; j++) {
+				tk::dnn::writeBUF(buf, tmp[j]);
+			}
+		}
+	}
+
+	int c, h, w;
+    int classes, num;
+	std::vector<std::string> classesNames;
+
+    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;
+	}
+
+};