Merge with master, all tests passed

Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>

include/tkDNN/pluginsRT/ActivationLeakyRT.h

@@ -51,9 +51,9 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
-		tk::dnn::writeBUF(buf, slope);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
 		tk::dnn::writeBUF(buf, size);
+		assert(buf == a + getSerializationSize());
 	}
 
 	int size;

include/tkDNN/pluginsRT/ActivationLogisticRT.h

@@ -0,0 +1,60 @@
+#include<cassert>
+#include "../kernels.h"
+
+class ActivationLogisticRT : public IPlugin {
+
+public:
+    ActivationLogisticRT() {
+
+
+    }
+
+    ~ActivationLogisticRT(){
+
+    }
+
+    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 {
+
+        activationLOGISTICForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
+                                  reinterpret_cast<dnnType*>(outputs[0]), batchSize*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/ActivationMishRT.h

@@ -52,8 +52,9 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
 		tk::dnn::writeBUF(buf, size);
+		assert(buf == a + getSerializationSize());
 	}
 
 	int size;

include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h

@@ -51,9 +51,10 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
 		tk::dnn::writeBUF(buf, ceiling);
 		tk::dnn::writeBUF(buf, size);
+		assert(buf = a + getSerializationSize());
 		
 	}
 

include/tkDNN/pluginsRT/ActivationSigmoidRT.h

@@ -52,8 +52,9 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
 		tk::dnn::writeBUF(buf, size);
+		assert(buf == a + getSerializationSize());
 	}
 
 	int size;

include/tkDNN/pluginsRT/DeformableConvRT.h

@@ -89,7 +89,7 @@ public:
 		for(int b=0; b<batchSize; b++) {
 			checkCuda(cudaMemcpy(offset, output_conv + b * 3 * chunk_dim, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); 
 			checkCuda(cudaMemcpy(mask, output_conv + b * 3 * chunk_dim + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); 
-			// kernel sigmoide
+			// kernel sigmoid
 			activationSIGMOIDForward(mask, mask, chunk_dim);
 			// deformable convolution
 			dcnV2CudaForward(stat, handle, 
@@ -116,7 +116,7 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
 		tk::dnn::writeBUF(buf, chunk_dim);
 		tk::dnn::writeBUF(buf, kh);
 		tk::dnn::writeBUF(buf, kw);
@@ -163,6 +163,7 @@ public:
         for(int i=0; i<dim_ones; i++)
     		tk::dnn::writeBUF(buf, aus[i]);
 		free(aus);
+		assert(buf == a + getSerializationSize());
 	}
 
 	cublasStatus_t stat; 

include/tkDNN/pluginsRT/FlattenConcatRT.h

@@ -65,12 +65,13 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a = buf;
 		tk::dnn::writeBUF(buf, c);
 		tk::dnn::writeBUF(buf, h);
 		tk::dnn::writeBUF(buf, w);
 		tk::dnn::writeBUF(buf, rows);
 		tk::dnn::writeBUF(buf, cols);
+		assert(buf == a + getSerializationSize());
 	}
 
 	int c, h, w;

include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h

@@ -55,7 +55,7 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
 
 		tk::dnn::writeBUF(buf, this->c);
 		tk::dnn::writeBUF(buf, this->h);
@@ -65,6 +65,7 @@ public:
 		tk::dnn::writeBUF(buf, this->stride_W);
 		tk::dnn::writeBUF(buf, this->winSize);
 		tk::dnn::writeBUF(buf, this->padding);
+		assert(buf == a + getSerializationSize());
 	}
 
 	int n, c, h, w;

include/tkDNN/pluginsRT/RegionRT.h

@@ -73,13 +73,14 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
 		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);
+		assert(buf == a + getSerializationSize());
 	}
 
 	int c, h, w;

include/tkDNN/pluginsRT/ReorgRT.h

@@ -52,11 +52,12 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
 		tk::dnn::writeBUF(buf, stride);
 		tk::dnn::writeBUF(buf, c);
 		tk::dnn::writeBUF(buf, h);
 		tk::dnn::writeBUF(buf, w);
+		assert(buf == a + getSerializationSize());
 	}
 
 	int c, h, w, stride;

include/tkDNN/pluginsRT/ReshapeRT.h

@@ -50,11 +50,12 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a = buf;
 		tk::dnn::writeBUF(buf, n);
 		tk::dnn::writeBUF(buf, c);
 		tk::dnn::writeBUF(buf, h);
 		tk::dnn::writeBUF(buf, w);
+		assert(buf == a + getSerializationSize());
 	}
 
 	int n, c, h, w;

include/tkDNN/pluginsRT/ResizeLayerRT.h

@@ -52,7 +52,7 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
 
 		tk::dnn::writeBUF(buf, o_c);
 		tk::dnn::writeBUF(buf, o_h);
@@ -61,6 +61,7 @@ public:
 		tk::dnn::writeBUF(buf, i_c);
 		tk::dnn::writeBUF(buf, i_h);
 		tk::dnn::writeBUF(buf, i_w);
+		assert(buf == a + getSerializationSize());
 	}
 
 	int i_c, i_h, i_w, o_c, o_h, o_w;

include/tkDNN/pluginsRT/RouteRT.h

@@ -8,7 +8,9 @@ class RouteRT : public IPlugin {
 	*/
 
 public:
-	RouteRT() {
+	RouteRT(int groups, int group_id) {
+		this->groups = groups;
+		this->group_id = group_id;
 	}
 
 	~RouteRT(){
@@ -22,7 +24,7 @@ public:
 	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]};
+		return DimsCHW{out_c/groups, inputs[0].d[1], inputs[0].d[2]};
 	}
 
 	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
@@ -34,6 +36,7 @@ public:
 		}
 		h = inputDims[0].d[1];
 		w = inputDims[0].d[2];
+		c /= groups;
 	}
 
 	int initialize() override {
@@ -49,15 +52,18 @@ public:
 	}
 
 	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;
+		for(int b=0; b<batchSize; b++) {
+			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;
+				int part_in_dim = in_dim / this->groups;
+				checkCuda( cudaMemcpyAsync(dstData + b*c*w*h + offset, input + b*c*w*h*groups + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) );
+				offset += part_in_dim;
+			}
 		}
 
 		return 0;
@@ -65,11 +71,13 @@ public:
 
 
 	virtual size_t getSerializationSize() override {
-		return (4+MAX_INPUTS)*sizeof(int);
+		return (6+MAX_INPUTS)*sizeof(int);
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
+		tk::dnn::writeBUF(buf, groups);
+		tk::dnn::writeBUF(buf, group_id);
 		tk::dnn::writeBUF(buf, in);
 		for(int i=0; i<MAX_INPUTS; i++)
 			tk::dnn::writeBUF(buf, c_in[i]);
@@ -77,10 +85,12 @@ public:
 		tk::dnn::writeBUF(buf, c);
 		tk::dnn::writeBUF(buf, h);
 		tk::dnn::writeBUF(buf, w);
+		assert(buf == a + getSerializationSize());
 	}
 
 	static const int MAX_INPUTS = 4;
 	int in;
 	int c_in[MAX_INPUTS];
 	int c, h, w;
+	int groups, group_id;
 };

include/tkDNN/pluginsRT/ShortcutRT.h

@@ -59,7 +59,7 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
 		tk::dnn::writeBUF(buf, bc);
 		tk::dnn::writeBUF(buf, bh);
 		tk::dnn::writeBUF(buf, bw);
@@ -67,7 +67,8 @@ public:
 		tk::dnn::writeBUF(buf, c);
 		tk::dnn::writeBUF(buf, h);
 		tk::dnn::writeBUF(buf, w);
-			
+		assert(buf == a + getSerializationSize());
+		
 	}
 
 	int c, h, w;

include/tkDNN/pluginsRT/UpsampleRT.h

@@ -54,11 +54,12 @@ public:
 	}
 
 	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
 		tk::dnn::writeBUF(buf, stride);
 		tk::dnn::writeBUF(buf, c);
 		tk::dnn::writeBUF(buf, h);
 		tk::dnn::writeBUF(buf, w);
+		assert(buf == a + getSerializationSize());
 	}
 
 	int c, h, w, stride;

include/tkDNN/pluginsRT/YoloRT.h

@@ -8,12 +8,15 @@ class YoloRT : public IPlugin {
 
 
 public:
-	YoloRT(int classes, int num, tk::dnn::Yolo *yolo = nullptr, int n_masks=3, float scale_xy=1) {
+	YoloRT(int classes, int num, tk::dnn::Yolo *yolo = nullptr, int n_masks=3, float scale_xy=1, float nms_thresh=0.45, int nms_kind=0, int new_coords=0) {
 
 		this->classes = classes;
 		this->num = num;
 		this->n_masks = n_masks;
 		this->scaleXY = scale_xy;
+		this->nms_thresh = nms_thresh;
+		this->nms_kind = nms_kind;
+		this->new_coords = new_coords;
 
         mask = new dnnType[n_masks];
         bias = new dnnType[num*n_masks*2];
@@ -61,17 +64,23 @@ public:
 
 		checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
 
-		for (int b = 0; b < batchSize; ++b){
-			for(int n = 0; n < n_masks; ++n){
-				int index = entry_index(b, n*w*h, 0);
-				activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream);
 
-				if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
-				
-				index = entry_index(b, n*w*h, 4);
-				activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream);
-			}
-		}
+        for (int b = 0; b < batchSize; ++b){
+            for(int n = 0; n < n_masks; ++n){
+                int index = entry_index(b, n*w*h, 0);
+                if (new_coords == 1){
+                    if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
+                }
+                else{
+                    activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); //x,y
+
+                    if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
+
+                    index = entry_index(b, n*w*h, 4);
+                    activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream);
+                }
+            }
+        }
 
 		//std::cout<<"YOLO END\n";
 		return 0;
@@ -79,22 +88,29 @@ public:
 
 
 	virtual size_t getSerializationSize() override {
-		return 6*sizeof(int) + sizeof(float)+ n_masks*sizeof(dnnType) + num*n_masks*2*sizeof(dnnType) + YOLORT_CLASSNAME_W*classes*sizeof(char);
+		return 8*sizeof(int) + 2*sizeof(float)+ n_masks*sizeof(dnnType) + num*n_masks*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, n_masks);
-		tk::dnn::writeBUF(buf, c);
-		tk::dnn::writeBUF(buf, h);
-		tk::dnn::writeBUF(buf, w);
-		tk::dnn::writeBUF(buf, scaleXY);
-        for(int i=0; i<n_masks; i++)
-    		tk::dnn::writeBUF(buf, mask[i]);
-        for(int i=0; i<n_masks*2*num; i++)
-    		tk::dnn::writeBUF(buf, bias[i]);
+		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
+		tk::dnn::writeBUF(buf, classes); 	//std::cout << "Classes :" << classes << std::endl;
+		tk::dnn::writeBUF(buf, num); 		//std::cout << "Num : " << num << std::endl;
+		tk::dnn::writeBUF(buf, n_masks); 	//std::cout << "N_Masks" << n_masks << std::endl;
+		tk::dnn::writeBUF(buf, scaleXY); 	//std::cout << "ScaleXY :" << scaleXY << std::endl;
+		tk::dnn::writeBUF(buf, nms_thresh); //std::cout << "nms_thresh :" << nms_thresh << std::endl;
+		tk::dnn::writeBUF(buf, nms_kind); 	//std::cout << "nms_kind : " << nms_kind << std::endl;
+		tk::dnn::writeBUF(buf, new_coords); //std::cout << "new_coords : " << new_coords << std::endl;
+		tk::dnn::writeBUF(buf, c); 			//std::cout << "C : " << c << std::endl;
+		tk::dnn::writeBUF(buf, h); 			//std::cout << "H : " << h << std::endl;
+		tk::dnn::writeBUF(buf, w); 			//std::cout << "C : " << c << std::endl;
+		for (int i = 0; i < n_masks; i++)
+		{
+			tk::dnn::writeBUF(buf, mask[i]); //std::cout << "mask[i] : " << mask[i] << std::endl;
+		}
+		for (int i = 0; i < n_masks * 2 * num; i++)
+		{
+			tk::dnn::writeBUF(buf, bias[i]); //std::cout << "bias[i] : " << bias[i] << std::endl;
+		}
 
 		// save classes names
 		for(int i=0; i<classes; i++) {
@@ -104,11 +120,15 @@ public:
 				tk::dnn::writeBUF(buf, tmp[j]);
 			}
 		}
+		assert(buf == a + getSerializationSize());
 	}
 
 	int c, h, w;
     int classes, num, n_masks;
 	float scaleXY;
+	float nms_thresh;
+	int nms_kind;
+	int new_coords;
 	std::vector<std::string> classesNames;
 
     dnnType *mask;