From 1a50065bdeed24ffd4956eb8a1f231d5b4461476 Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Tue, 17 Mar 2020 13:18:05 +0100 Subject: [PATCH] Fix conv2d with additional bias for tensorRT. Fix reshape deserialize. Mb2512 works with tensorRT Signed-off-by: Micaela Verucchi --- include/tkDNN/Layer.h | 1 + src/LayerWgs.cpp | 8 ++++++++ src/NetworkRT.cpp | 18 ++++++++++++++---- tests/mobilenetv2ssd512/mobilenetv2ssd512.cpp | 2 -- 4 files changed, 23 insertions(+), 6 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 0d6d472..625e945 100644 --- a/include/tkDNN/Layer.h +++ b/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; diff --git a/src/LayerWgs.cpp b/src/LayerWgs.cpp index 9dad5e9..ed1b62c 100644 --- a/src/LayerWgs.cpp +++ b/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)); diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 17f32bf..29b183d 100644 --- a/src/NetworkRT.cpp +++ b/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 - b = { dtRT, nullptr, 0}; //on batchnorm bias are added later + else{ + 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(buf), readBUF(buf),readBUF(buf), readBUF(buf)); + dataDim_t new_dim; + new_dim.n = readBUF(buf); + new_dim.c = readBUF(buf); + new_dim.h = readBUF(buf); + new_dim.w = readBUF(buf); ReshapeRT *r = new ReshapeRT(new_dim); return r; diff --git a/tests/mobilenetv2ssd512/mobilenetv2ssd512.cpp b/tests/mobilenetv2ssd512/mobilenetv2ssd512.cpp index eb2887e..b207cec 100644 --- a/tests/mobilenetv2ssd512/mobilenetv2ssd512.cpp +++ b/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