Fix conv2d with additional bias for tensorRT. Fix reshape deserialize. Mb2512 works with tensorRT

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

include/tkDNN/Layer.h

@@ -114,6 +114,7 @@ public:
     //fp16
     __half *data16_h, *bias16_h;
     __half *data16_d, *bias16_d;
+    __half *bias216_h, *bias216_d;
 
     __half *power16_h,    *power16_d;
     __half *scales16_h,   *scales16_d;

src/LayerWgs.cpp

@@ -59,6 +59,14 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs,
     float2half(data_d, data16_d, w_size);
     cudaMemcpy(data16_h, data16_d, w_size*sizeof(__half), cudaMemcpyDeviceToHost);
 
+    if(additional_bias){
+        int b2_size = outputs;
+        bias216_h = new __half[b2_size];
+        cudaMalloc(&bias216_d, w_size*sizeof(__half));
+        float2half(bias2_d, bias216_d, b2_size);
+        cudaMemcpy(bias216_h, bias216_d, b2_size*sizeof(__half), cudaMemcpyDeviceToHost);
+    }
+
     int b_size = outputs;
     bias16_h = new __half[b_size];
     cudaMalloc(&bias16_d, w_size*sizeof(__half));

src/NetworkRT.cpp

@@ -226,10 +226,11 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
     // printf("%d %d %d %d %d\n", l->kernelH, l->kernelW, l->inputs, l->outputs, l->batchnorm);
 
 
-    void *data_b, *bias_b, *power_b, *mean_b, *variance_b, *scales_b;
+    void *data_b, *bias_b, *bias2_b, *power_b, *mean_b, *variance_b, *scales_b;
     if(dtRT == DataType::kHALF) {
         data_b     = l->data16_h;    
         bias_b     = l->bias16_h;
+        bias2_b    = l->bias216_h;
         power_b    = l->power16_h;
         mean_b     = l->mean16_h;
         variance_b = l->variance16_h;
@@ -237,6 +238,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
     } else {
         data_b     = l->data_h;    
         bias_b     = l->bias_h;
+        bias2_b    = l->bias2_h;
         power_b    = l->power_h;
         mean_b     = l->mean_h;
         variance_b = l->variance_h;
@@ -248,8 +250,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
     Weights b;
     if(!l->batchnorm)
         b = { dtRT, bias_b, l->outputs};
-    else
+    else{
-        b = { dtRT, nullptr, 0}; //on batchnorm bias are added later
+        if (l->additional_bias)
+            b = { dtRT, bias2_b, l->outputs}; 
+        else
+            b = { dtRT, nullptr, 0}; //on batchnorm bias are added later
+    }
 
     ILayer *lRT = nullptr;
     if(!l->deConv) {
@@ -652,7 +658,11 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
 
     if(name.find("Reshape") == 0) {
 
-        dataDim_t new_dim(readBUF<int>(buf), readBUF<int>(buf),readBUF<int>(buf), readBUF<int>(buf));
+        dataDim_t new_dim;
+        new_dim.n = readBUF<int>(buf);
+        new_dim.c = readBUF<int>(buf);
+        new_dim.h = readBUF<int>(buf);
+        new_dim.w = readBUF<int>(buf);
         ReshapeRT *r = new ReshapeRT(new_dim); 
         
         return r;

tests/mobilenetv2ssd512/mobilenetv2ssd512.cpp

@@ -445,8 +445,6 @@ int main()
     tk::dnn::Reshape reshape_conf2(&net, newdim_c);
 
     tk::dnn::Softmax sm_1(&net, &newdim_c, true);
-    // tk::dnn::Flatten fl_l_7(&net);
-    // tk::dnn::Reshape reshape_conf3(&net,dim_resh, true);
     tk::dnn::Layer *conf = &sm_1;
 
     //concat locations