From 5272f1cde6c5ac06a786cafac4d04102ea2971ff Mon Sep 17 00:00:00 2001 From: Davide Sapienza Date: Mon, 23 Dec 2019 15:31:13 +0100 Subject: [PATCH] CenterNet TensorRT serialization works This commit adds the Deformable layer serialization. Signed-oof-by: Davide Sapienza --- include/tkDNN/Layer.h | 2 + include/tkDNN/pluginsRT/DeformableConvRT.h | 153 ++++++++++++++++++--- src/Network.cpp | 1 + src/NetworkRT.cpp | 56 +++++++- 4 files changed, 188 insertions(+), 24 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index a6c639b..f5ac6a7 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -43,6 +43,8 @@ public: dataDim_t input_dim, output_dim; dnnType *dstData; //where results will be putted + + int id = 0; bool final; //if the layer is the final one std::string getLayerName() { diff --git a/include/tkDNN/pluginsRT/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h index 6a85538..b236095 100644 --- a/include/tkDNN/pluginsRT/DeformableConvRT.h +++ b/include/tkDNN/pluginsRT/DeformableConvRT.h @@ -7,8 +7,51 @@ class DeformableConvRT : public IPlugin { public: - DeformableConvRT(tk::dnn::DeformConv2d *deformable) { - this->defRT = deformable; + DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw, + int deformableGroup, int i_n, int i_c, int i_h, int i_w, + int o_n, int o_c, int o_h, int o_w, + tk::dnn::DeformConv2d *deformable = nullptr) { + this->chunk_dim = chunk_dim; + // int dst_dim = conv_dim.tot(); + // std::cout<<"conv_dim: \n"; + // conv_dim.print(); + // if (dst_dim % 3 != 0 ) + // std::cout<<"take attention\n\n"; + // this->chunk_dim = dst_dim/3; + this->kh = kh; + this->kw = kw; + this->sh = sh; + this->sw = sw; + this->ph = ph; + this->pw = pw; + this->deformableGroup = deformableGroup; + this->i_n = i_n; + this->i_c = i_c; + this->i_h = i_h; + this->i_w = i_w; + this->o_n = o_n; + this->o_c = o_c; + this->o_h = o_h; + this->o_w = o_w; + height_ones = (i_h + 2 * ph - (1 * (kh - 1) + 1)) / sh + 1; + width_ones = (i_w + 2 * pw - (1 * (kw - 1) + 1)) / sw + 1; + dim_ones = i_c * kh * kw * 1 * height_ones * width_ones; + std::cout<defRT = deformable; + checkCuda( cudaMemcpy(data_d, deformable->data_d, sizeof(dnnType)*i_c * o_c * kh * kw * 1, cudaMemcpyDeviceToDevice) ); + checkCuda( cudaMemcpy(bias2_d, deformable->bias2_d, sizeof(dnnType)*o_c, cudaMemcpyDeviceToDevice) ); + checkCuda( cudaMemcpy(ones_d1, deformable->ones_d1, sizeof(dnnType)*height_ones*width_ones, cudaMemcpyDeviceToDevice) ); + checkCuda( cudaMemcpy(offset, deformable->offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToDevice) ); + checkCuda( cudaMemcpy(mask, deformable->mask, sizeof(dnnType)*chunk_dim, cudaMemcpyDeviceToDevice) ); + checkCuda( cudaMemcpy(ones_d2, deformable->ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToDevice) ); + } } ~DeformableConvRT(){ @@ -24,6 +67,14 @@ public: } void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + // i_n = 1; + // i_c = inputDims[0].d[0]; + // i_h = inputDims[0].d[1]; + // i_w = inputDims[0].d[2]; + // o_n = 1; + // o_c = outputDims[0].d[0]; + // o_h = outputDims[0].d[1]; + // o_w = outputDims[0].d[2]; } int initialize() override { @@ -39,42 +90,108 @@ public: } virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { - +std::cout<<"LOL\n"; dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); dnnType *output_conv = (dnnType*)reinterpret_cast(inputs[1]); // split conv2d outputs into offset to mask - checkCuda(cudaMemcpy(defRT->offset, defRT->output_conv, 2*defRT->chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); - checkCuda(cudaMemcpy(defRT->mask, defRT->output_conv + 2*defRT->chunk_dim, defRT->chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + checkCuda(cudaMemcpy(offset, output_conv, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + checkCuda(cudaMemcpy(mask, output_conv + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); // kernel sigmoide - activationSIGMOIDForward(defRT->mask, defRT->mask, defRT->chunk_dim); + activationSIGMOIDForward(mask, mask, chunk_dim); // deformable convolution - dcn_v2_cuda_forward(srcData, defRT->data_d, - defRT->bias2_d, defRT->ones_d1, - defRT->offset, defRT->mask, - reinterpret_cast(outputs[0]), defRT->ones_d2, - defRT->kernelH, defRT->kernelW, - defRT->strideH, defRT->strideW, - defRT->paddingH, defRT->paddingW, + dcn_v2_cuda_forward(srcData, data_d, + bias2_d, ones_d1, + offset, mask, + reinterpret_cast(outputs[0]), ones_d2, + kh, kw, + sh, sw, + ph, pw, 1, 1, - defRT->deformableGroup, - defRT->preconv->input_dim.n, defRT->preconv->input_dim.c, defRT->preconv->input_dim.h, defRT->preconv->input_dim.w, - defRT->output_dim.n, defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w, - defRT->chunk_dim); + deformableGroup, + i_n, i_c, i_h, i_w, + o_n, o_c, o_h, o_w, + chunk_dim); return 0; } virtual size_t getSerializationSize() override { - return 0; + return 16 * sizeof(int) + chunk_dim * 3 * sizeof(dnnType) + (i_c * o_c * kh * kw * 1 ) * sizeof(dnnType) + + o_c * sizeof(dnnType) + height_ones * width_ones * sizeof(dnnType) + dim_ones * sizeof(dnnType); } virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer); + tk::dnn::writeBUF(buf, chunk_dim); + tk::dnn::writeBUF(buf, kh); + tk::dnn::writeBUF(buf, kw); + tk::dnn::writeBUF(buf, sh); + tk::dnn::writeBUF(buf, sw); + tk::dnn::writeBUF(buf, ph); + tk::dnn::writeBUF(buf, pw); + tk::dnn::writeBUF(buf, deformableGroup); + tk::dnn::writeBUF(buf, i_n); + tk::dnn::writeBUF(buf, i_c); + tk::dnn::writeBUF(buf, i_h); + tk::dnn::writeBUF(buf, i_w); + tk::dnn::writeBUF(buf, o_n); + tk::dnn::writeBUF(buf, o_c); + tk::dnn::writeBUF(buf, o_h); + tk::dnn::writeBUF(buf, o_w); + dnnType *aus = new dnnType[chunk_dim*2]; + checkCuda( cudaMemcpy(aus, offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToHost) ); + for(int i=0; iid = num_layers; layers[num_layers++] = l; return true; } diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index ed72c89..117121d 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -426,17 +426,21 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Upsample *l) { } ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) { - //std::cout<<"convert DEFORMABLE\n"; + std::cout<<"convert DEFORMABLE\n"; ILayer *preconv = convert_layer(input, l->preconv); + checkNULL(preconv); ITensor **inputs = new ITensor*[2]; inputs[0] = input; inputs[1] = preconv->getOutput(0); - //std::cout<<"New plugin DEFORMABLE\n"; - IPlugin *plugin = new DeformableConvRT(l); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + std::cout<<"New plugin DEFORMABLE\n"; + IPlugin *plugin = new DeformableConvRT(l->chunk_dim, l->kernelH, l->kernelW, l->strideH, l->strideW, l->paddingH, l->paddingW, + l->deformableGroup, l->input_dim.n, l->input_dim.c, l->input_dim.h, l->input_dim.w, + l->output_dim.n, l->output_dim.c, l->output_dim.h, l->output_dim.w, l); + IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin); checkNULL(lRT); + lRT->setName( ("Deformable" + std::to_string(l->id)).c_str() ); // batchnorm void *bias_b, *power_b, *mean_b, *variance_b, *scales_b; @@ -515,8 +519,9 @@ bool NetworkRT::deserialize(const char *filename) { IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) { const char * buf = reinterpret_cast(serialData); - + std::string name(layerName); + std::cout<(buf)); //stride r->c = readBUF(buf); @@ -606,6 +610,46 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa return r; } */ + if(name.find("Deformable") == 0) { + DeformableConvRT *r = new DeformableConvRT(readBUF(buf), readBUF(buf), readBUF(buf), + readBUF(buf), readBUF(buf), readBUF(buf), + readBUF(buf), readBUF(buf), + readBUF(buf),readBUF(buf),readBUF(buf),readBUF(buf), + readBUF(buf),readBUF(buf),readBUF(buf),readBUF(buf), + nullptr); + dnnType *aus = new dnnType[r->chunk_dim*2]; + for(int i=0; ichunk_dim*2; i++) + aus[i] = readBUF(buf); + checkCuda( cudaMemcpy(r->offset, aus, sizeof(dnnType)*2*r->chunk_dim, cudaMemcpyHostToDevice) ); + free(aus); + aus = new dnnType[r->chunk_dim]; + for(int i=0; ichunk_dim; i++) + aus[i] = readBUF(buf); + checkCuda( cudaMemcpy(r->mask, aus, sizeof(dnnType)*r->chunk_dim, cudaMemcpyHostToDevice) ); + free(aus); + aus = new dnnType[(r->i_c * r->o_c * r->kh * r->kw * 1 )]; + for(int i=0; i<(r->i_c * r->o_c * r->kh * r->kw * 1 ); i++) + aus[i] = readBUF(buf); + checkCuda( cudaMemcpy(r->data_d, aus, sizeof(dnnType)*(r->i_c * r->o_c * r->kh * r->kw * 1 ), cudaMemcpyHostToDevice) ); + free(aus); + aus = new dnnType[r->o_c]; + for(int i=0; i < r->o_c; i++) + aus[i] = readBUF(buf); + checkCuda( cudaMemcpy(r->bias2_d, aus, sizeof(dnnType)*r->o_c, cudaMemcpyHostToDevice) ); + free(aus); + aus = new dnnType[r->height_ones * r->width_ones]; + for(int i=0; iheight_ones * r->width_ones; i++) + aus[i] = readBUF(buf); + checkCuda( cudaMemcpy(r->ones_d1, aus, sizeof(dnnType)*r->height_ones * r->width_ones, cudaMemcpyHostToDevice) ); + free(aus); + aus = new dnnType[r->dim_ones]; + for(int i=0; idim_ones; i++) + aus[i] = readBUF(buf); + checkCuda( cudaMemcpy(r->ones_d2, aus, sizeof(dnnType)*r->dim_ones, cudaMemcpyHostToDevice) ); + free(aus); + return r; + } + FatalError("Cant deserialize Plugin"); return NULL; }