From 9f1e30eaa9c637e9ef30ab78e9f4809f4caac4e2 Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Fri, 19 Jun 2020 18:55:42 +0200 Subject: [PATCH 01/17] Add shelfnet. Resnet18backbone works Signed-off-by: Micaela Verucchi --- CMakeLists.txt | 4 + src/kernels/activation_leaky.cu | 2 +- src/utils.cpp | 1 + tests/shelfnet/shelfnet.cpp | 213 ++++++++++++++++++++++++++++++++ 4 files changed, 219 insertions(+), 1 deletion(-) create mode 100644 tests/shelfnet/shelfnet.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 8c8619d..a2a8a06 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -103,6 +103,10 @@ target_link_libraries(test_resnet101_cnet tkDNN) add_executable(test_dla34_cnet tests/centernet/dla34_cnet/dla34_cnet.cpp) target_link_libraries(test_dla34_cnet tkDNN) +# SHELFNET +add_executable(test_shelfnet tests/shelfnet/shelfnet.cpp) +target_link_libraries(test_shelfnet tkDNN) + # DEMOS add_executable(test_rtinference tests/test_rtinference/rtinference.cpp) target_link_libraries(test_rtinference tkDNN) diff --git a/src/kernels/activation_leaky.cu b/src/kernels/activation_leaky.cu index a9029ad..47dc18f 100644 --- a/src/kernels/activation_leaky.cu +++ b/src/kernels/activation_leaky.cu @@ -9,7 +9,7 @@ void activation_leaky(dnnType *input, dnnType *output, int size) { if (input[i]>0) output[i] = input[i]; else - output[i] = 0.1f*input[i]; + output[i] = 0.01f*input[i]; //FIME!! } } diff --git a/src/utils.cpp b/src/utils.cpp index 65030f0..87aad06 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -102,6 +102,7 @@ int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device, int } int diffs = 0; for(int i=0; i eps) { diffs += 1; diff --git a/tests/shelfnet/shelfnet.cpp b/tests/shelfnet/shelfnet.cpp new file mode 100644 index 0000000..3f8834a --- /dev/null +++ b/tests/shelfnet/shelfnet.cpp @@ -0,0 +1,213 @@ +#include +#include "tkdnn.h" + + +const char *output_bin1 = "shelfnet/debug/classification_headers-5.bin"; +const char *output_bin2 = "shelfnet/debug/regression_headers-5.bin"; +const char *input_bin = "shelfnet/debug/input.bin"; + +const char *backbone[] = { + "shelfnet/layers/backbone-conv1.bin", + "shelfnet/layers/backbone-layer1-0-conv1.bin", + "shelfnet/layers/backbone-layer1-0-conv2.bin", + "shelfnet/layers/backbone-layer1-1-conv1.bin", + "shelfnet/layers/backbone-layer1-1-conv2.bin", + "shelfnet/layers/backbone-layer2-0-conv1.bin", + "shelfnet/layers/backbone-layer2-0-conv2.bin", + "shelfnet/layers/backbone-layer2-0-downsample-0.bin", + "shelfnet/layers/backbone-layer2-1-conv1.bin", + "shelfnet/layers/backbone-layer2-1-conv2.bin", + "shelfnet/layers/backbone-layer3-0-conv1.bin", + "shelfnet/layers/backbone-layer3-0-conv2.bin", + "shelfnet/layers/backbone-layer3-0-downsample-0.bin", + "shelfnet/layers/backbone-layer3-1-conv1.bin", + "shelfnet/layers/backbone-layer3-1-conv2.bin", + "shelfnet/layers/backbone-layer4-0-conv1.bin", + "shelfnet/layers/backbone-layer4-0-conv2.bin", + "shelfnet/layers/backbone-layer4-0-downsample-0.bin", + "shelfnet/layers/backbone-layer4-1-conv1.bin", + "shelfnet/layers/backbone-layer4-1-conv2.bin"}; + +const char *conv_out[] = { + "shelfnet/layers/conv_out16-conv-conv.bin", + "shelfnet/layers/conv_out16-conv_out.bin", + "shelfnet/layers/conv_out32-conv-conv.bin", + "shelfnet/layers/conv_out32-conv_out.bin", + "shelfnet/layers/conv_out-conv-conv.bin", + "shelfnet/layers/conv_out-conv_out.bin"}; + +const char *decoder[] = { + "shelfnet/layers/decoder-bottom-conv1.bin", + "shelfnet/layers/decoder-up_conv_list-0-conv_atten.bin", + "shelfnet/layers/decoder-up_conv_list-0-conv-conv.bin", + "shelfnet/layers/decoder-up_conv_list-1-conv_atten.bin", + "shelfnet/layers/decoder-up_conv_list-1-conv-conv.bin", + "shelfnet/layers/decoder-up_dense_list-0-conv.bin", + "shelfnet/layers/decoder-up_dense_list-1-conv.bin"}; + + +const char *ladder[] = { + "shelfnet/layers/ladder-bottom-conv1.bin", + "shelfnet/layers/ladder-down_conv_list-0.bin", + "shelfnet/layers/ladder-down_conv_list-1.bin", + "shelfnet/layers/ladder-down_module_list-0-conv1.bin", + "shelfnet/layers/ladder-down_module_list-1-conv1.bin", + "shelfnet/layers/ladder-inconv-conv1.bin", + "shelfnet/layers/ladder-up_conv_list-0-conv_atten.bin", + "shelfnet/layers/ladder-up_conv_list-0-conv-conv.bin", + "shelfnet/layers/ladder-up_conv_list-1-conv_atten.bin", + "shelfnet/layers/ladder-up_conv_list-1-conv-conv.bin", + "shelfnet/layers/ladder-up_dense_list-0-conv.bin", + "shelfnet/layers/ladder-up_dense_list-1-conv.bin"}; + +const char *trans[] = { + "shelfnet/layers/trans1-conv.bin", + "shelfnet/layers/trans2-conv.bin", + "shelfnet/layers/trans3-conv.bin"}; +int main() +{ + + // downloadWeightsifDoNotExist(input_bin, "shelfnet", "https://cloud.hipert.unimore.it/s/x4ZfxBKN23zAJQp/download"); + + int classes = 19; + + // Network layout + tk::dnn::dataDim_t dim(1, 3, 1024, 1024, 1); + tk::dnn::Network net(dim); + + int bi = 0; + new tk::dnn::Conv2d(&net, 64, 7, 7, 2, 2, 3, 3, backbone[bi++], true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Layer* last = new tk::dnn::Pooling (&net, 3, 3, 2, 2, 1, 1, tk::dnn::POOLING_MAX); + + + + for(int i=0; i<2; ++i){ + new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, backbone[bi++], true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY); + new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, backbone[bi++], true); + new tk::dnn::Shortcut(&net, last); + last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); + } + + std::vector features; + for(int i=0;i<3;++i){ + int out_channel = pow(2,7+i); + std::cout< batch_frame; + std::vector batch_dnn_input; + + while(gRun) { + batch_dnn_input.clear(); + batch_frame.clear(); + + for(int bi=0; bi< n_batch; ++bi){ + cap >> frame; + if(!frame.data) + break; + + batch_frame.push_back(frame); + + // this will be resized to the net format + batch_dnn_input.push_back(frame.clone()); + } + if(!frame.data) + break; + + //inference + segNN.update(batch_dnn_input, n_batch); + segNN.draw(); + + if(show){ + for(int bi=0; bi< n_batch; ++bi){ + cv::imshow("segmentation", batch_frame[bi]); + cv::waitKey(1); + } + } + if(n_batch == 1 && SAVE_RESULT) + resultVideo << frame; + } + + std::cout<<"segmentation end\n"; + double mean = 0; + + std::cout< +#include +#include +#include +#include +#include "utils.h" + +#include +#include +#include +#include + + + +#include "tkdnn.h" +#include "NetworkViz.h" + +namespace tk { namespace dnn { + +class SegmentationNN { + + protected: + tk::dnn::NetworkRT *netRT = nullptr; + int nBatches = 1; + + std::vector originalSize; + std::vector masks; + cv::Mat bgr[3]; + dnnType *input; + dnnType *input_d; + float* confidences_h; + + /** + * This method preprocess the image, before feeding it to the NN. + * + * @param frame original frame to adapt for inference. + * @param bi batch index + */ + void preprocess(cv::Mat &frame, const int bi=0) { + + frame.convertTo(frame, CV_32FC3, 1 / 255.0, 0); + + cv::split(frame, bgr); + float mean[] = {0.485, 0.456, 0.406}; + float stddev[] = {0.229, 0.224, 0.225}; + for(int i=0; i<3; i++){ + bgr[2-i] -= mean[i]; + bgr[2-i] /= stddev[i]; + } + cv::merge(bgr, 3, frame); + + int crop_size = netRT->input_dim.w; + int H = frame.rows; + int W = frame.cols; + cv::Mat frame_cropped; + + cv::Mat mask(frame.size(), CV_8UC3, cv::Scalar(255,255,255)); + + if(H != W){ + if(H < W){ + int top = (W - H)/2; + int bottom = W - top - H; + cv::copyMakeBorder(frame, frame_cropped, top, bottom, 0, 0, cv::BORDER_CONSTANT, cv::Scalar(0,0,0) ); + cv::copyMakeBorder(mask, mask, top, bottom, 0, 0, cv::BORDER_CONSTANT, cv::Scalar(0,0,0) ); + } + else{ + int left = (H - W)/2; + int right = H - left - W; + cv::copyMakeBorder(frame, frame_cropped, 0, 0, left, right, cv::BORDER_CONSTANT, cv::Scalar(0,0,0) ); + cv::copyMakeBorder(mask, mask, 0, 0, left, right, cv::BORDER_CONSTANT, cv::Scalar(0,0,0) ); + } + } + + resize(frame_cropped, frame_cropped, cv::Size(netRT->input_dim.w, netRT->input_dim.h)); + resize(mask, mask, cv::Size(netRT->input_dim.w, netRT->input_dim.h)); + masks[bi] = mask.clone(); + + cv::split(frame_cropped, bgr); + for (int i = 0; i < netRT->input_dim.c; i++){ + int idx = i * frame_cropped.rows * frame_cropped.cols; + int ch = netRT->input_dim.c-1 -i; + memcpy((void *)&input[idx + netRT->input_dim.tot()*bi], (void *)bgr[ch].data, frame_cropped.rows * frame_cropped.cols * sizeof(dnnType)); + } + checkCuda(cudaMemcpyAsync(input_d+ netRT->input_dim.tot()*bi, input + netRT->input_dim.tot()*bi, netRT->input_dim.tot() * sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream)); + } + + /** + * This method postprocess the output of the NN to obtain the correct + * boundig boxes. + * + * @param bi batch index + */ + void postprocess(const int bi=0) { + dnnType *rt_out = (dnnType *)netRT->buffersRT[1]+ netRT->buffersDIM[1].tot()*bi; + + dataDim_t odim = netRT->output_dim; + + + checkCuda(cudaMemcpy(confidences_h, rt_out, odim.tot() * sizeof(float), cudaMemcpyDeviceToHost)); + + for(int i=0;i max_conf){ + max_conf = cur_conf; + max_id = k; + } + } + confidences_h[bi*odim.tot()+0*odim.h*odim.w+i*odim.h+j] = max_id; + } + } + dataDim_t vdim = odim; + vdim.c = 1; + segmented[bi] = vizData2Mat(confidences_h, vdim, 1024, 0, 18); + }; + + public: + int classes = 0; + std::vector stats; /*keeps track of inference times (ms)*/ + std::vector classesNames; + std::vector segmented; + + SegmentationNN() {}; + ~SegmentationNN(){}; + + /** + * Method used to inialize the class, allocate memory and compute + * needed data. + * + * @param tensor_path path to the rt file og the NN. + * @param n_classes number of classes for the given dataset. + * @param n_batches maximum number of batches to use in inference + * @return true if everything is correct, false otherwise. + */ + bool init(const std::string& tensor_path, const int n_classes=19, const int n_batches=1){ + std::cout<<(tensor_path).c_str()<<"\n"; + if(!fileExist(tensor_path.c_str())) + FatalError("This file do not exists" + tensor_path ); + + netRT = new tk::dnn::NetworkRT(NULL, (tensor_path).c_str()); + classes = n_classes; + nBatches = n_batches; + + checkCuda(cudaMallocHost(&input, sizeof(dnnType) * netRT->input_dim.tot() * nBatches)); + checkCuda(cudaMalloc(&input_d, sizeof(dnnType) * netRT->input_dim.tot() * nBatches)); + + confidences_h = (float *)malloc(netRT->output_dim.tot() * sizeof(float)); + + segmented.resize(nBatches); + masks.resize(nBatches); + } + + /** + * This method performs the whole detection of the NN. + * + * @param frames frames to run detection on. + * @param cur_batches number of batches to use in inference + * @param save_times if set to true, preprocess, inference and postprocess times + * are saved on a csv file, otherwise not. + * @param times pointer to the output stream where to write times + * @param mAP set to true only if all the probabilities for a bounding + * box are needed, as in some cases for the mAP calculation + */ + void update(std::vector& frames, const int cur_batches=1){ + if(cur_batches > nBatches) + FatalError("A batch size greater than nBatches cannot be used"); + + originalSize.clear(); + if(TKDNN_VERBOSE) printCenteredTitle(" TENSORRT detection ", '=', 30); + { + TKDNN_TSTART + for(int bi=0; biinput_dim; + dim.n = cur_batches; + { + if(TKDNN_VERBOSE) dim.print(); + TKDNN_TSTART + netRT->infer(dim, input_d); + TKDNN_TSTOP + if(TKDNN_VERBOSE) dim.print(); + stats.push_back(t_ns); + } + + { + TKDNN_TSTART + for(int bi=0; bi> $out_file print_output $? imuodom + test_net shelfnet test_net yolo4 test_net yolo4_berkeley test_net yolo3 diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 71662d3..92a95f2 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -483,6 +483,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Resize *l) { IResizeLayer *lRT = networkRT->addResize(*input); //default is kNEAREST checkNULL(lRT); Dims d{}; + lRT->setResizeMode(ResizeMode(l->mode)); lRT->setOutputDimensions(DimsCHW{l->output_dim.c, l->output_dim.h, l->output_dim.w}); return lRT; } @@ -514,7 +515,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) { ITensor *back_tens = tensors[l->backLayer]; - if(false) //l->backLayer->output_dim.c == l->output_dim.c && !l->mul) FIXME + if(l->backLayer->output_dim.c == l->output_dim.c && !l->mul) { IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM); checkNULL(lRT); diff --git a/src/NetworkViz.cpp b/src/NetworkViz.cpp index 6ac274c..842a26e 100644 --- a/src/NetworkViz.cpp +++ b/src/NetworkViz.cpp @@ -6,23 +6,22 @@ namespace tk { namespace dnn { -cv::Mat vizFloat2colorMap(cv::Mat map) { +cv::Mat vizFloat2colorMap(cv::Mat map,double min, double max) { + + if(min == 0 && max == 0) + cv::minMaxIdx(map, &min, &max); - double min; - double max; - cv::minMaxIdx(map, &min, &max); cv::Mat adjMap; // expand your range to 0..255. Similar to histEq(); map.convertTo(adjMap,CV_8UC1, 255 / (max-min), -min); //return adjMap; - cv::Mat falseColorsMap; - applyColorMap(adjMap, falseColorsMap, cv::COLORMAP_HOT); + applyColorMap(adjMap, falseColorsMap, cv::COLORMAP_VIRIDIS); return falseColorsMap; } -cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim) { +cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim, double min, double max) { dnnType *data = nullptr; // copy to CPU @@ -38,7 +37,7 @@ cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim) { cv::Mat grid = cv::Mat(gridSize, CV_8UC3, cv::Scalar(0)); for(int i=0; imode = mode; if(fixed){ output_dim.c = scale_c; output_dim.h = scale_h; diff --git a/tests/shelfnet/shelfnet.cpp b/tests/shelfnet/shelfnet.cpp index ab28d90..7cc1b42 100644 --- a/tests/shelfnet/shelfnet.cpp +++ b/tests/shelfnet/shelfnet.cpp @@ -191,7 +191,7 @@ int main() down_out.push_back(l_last); new tk::dnn::Conv2d (&net, out_channel*2, 3, 3, 2, 2, 1, 1, ladder[li++], false); - last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); + last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.0f); //should be ReLU } new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); @@ -231,12 +231,12 @@ int main() new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, conv_out[ci++], true); new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); new tk::dnn::Conv2d (&net, 19, 3, 3, 1, 1, 1, 1, conv_out[ci++], false); - // /*up_out[i] =*/ new tk::dnn::Resize(&net, 19, net.input_dim.h, net.input_dim.w, true); + /*up_out[i] =*/ new tk::dnn::Resize(&net, 19, net.input_dim.h, net.input_dim.w, true, tk::dnn::ResizeMode_t::LINEAR); // } - // new tk::dnn::Softmax(&net); + new tk::dnn::Softmax(&net); - const char *output_bin = "shelfnet/debug/conv_out-conv_out.bin"; + const char *output_bin = "shelfnet/debug/softmax.bin"; // Load input dnnType *data; diff --git a/tests/test_rtinference/rtinference.cpp b/tests/test_rtinference/rtinference.cpp index a629168..3eacc37 100644 --- a/tests/test_rtinference/rtinference.cpp +++ b/tests/test_rtinference/rtinference.cpp @@ -31,7 +31,7 @@ int main(int argc, char *argv[]) { std::cout<<"Testing with batchsize: "< batch_frame; std::vector batch_dnn_input; @@ -85,14 +82,8 @@ int main(int argc, char *argv[]) { //inference segNN.update(batch_dnn_input, n_batch); - segNN.draw(); + frame = segNN.draw(); - if(show){ - for(int bi=0; bi< n_batch; ++bi){ - cv::imshow("segmentation", batch_frame[bi]); - cv::waitKey(1); - } - } if(n_batch == 1 && SAVE_RESULT) resultVideo << frame; } diff --git a/include/tkDNN/SegmentationNN.h b/include/tkDNN/SegmentationNN.h index 93be2ff..0289d61 100644 --- a/include/tkDNN/SegmentationNN.h +++ b/include/tkDNN/SegmentationNN.h @@ -17,6 +17,7 @@ #include "tkdnn.h" #include "NetworkViz.h" +#include "kernelsThrust.h" namespace tk { namespace dnn { @@ -33,6 +34,12 @@ class SegmentationNN { dnnType *input_d; float* confidences_h; + float * tmpInputData_d; + float *tmpOutData_d; + float *tmpOutData_h; + + cublasHandle_t cublasHandle; + /** * This method preprocess the image, before feeding it to the NN. * @@ -97,28 +104,14 @@ class SegmentationNN { dnnType *rt_out = (dnnType *)netRT->buffersRT[1]+ netRT->buffersDIM[1].tot()*bi; dataDim_t odim = netRT->output_dim; - - checkCuda(cudaMemcpy(confidences_h, rt_out, odim.tot() * sizeof(float), cudaMemcpyDeviceToHost)); + matrixTranspose(cublasHandle, rt_out, tmpInputData_d, odim.c, odim.w*odim.h); + maxElem(tmpInputData_d, tmpOutData_d, odim.c, odim.h, odim.w); + checkCuda(cudaMemcpy(tmpOutData_h, tmpOutData_d, odim.w*odim.h * sizeof(float), cudaMemcpyDeviceToHost)); - for(int i=0;i max_conf){ - max_conf = cur_conf; - max_id = k; - } - } - confidences_h[bi*odim.tot()+0*odim.h*odim.w+i*odim.h+j] = max_id; - } - } dataDim_t vdim = odim; vdim.c = 1; - segmented[bi] = vizData2Mat(confidences_h, vdim, 1024, 0, 18); + segmented[bi] = vizData2Mat(tmpOutData_h, vdim, 1024, 0, 18); }; public: @@ -127,8 +120,12 @@ class SegmentationNN { std::vector classesNames; std::vector segmented; - SegmentationNN() {}; - ~SegmentationNN(){}; + SegmentationNN() { + checkERROR( cublasCreate(&cublasHandle) ); + }; + ~SegmentationNN(){ + checkERROR( cublasDestroy(cublasHandle) ); + }; /** * Method used to inialize the class, allocate memory and compute @@ -151,7 +148,12 @@ class SegmentationNN { checkCuda(cudaMallocHost(&input, sizeof(dnnType) * netRT->input_dim.tot() * nBatches)); checkCuda(cudaMalloc(&input_d, sizeof(dnnType) * netRT->input_dim.tot() * nBatches)); - confidences_h = (float *)malloc(netRT->output_dim.tot() * sizeof(float)); + dataDim_t odim = netRT->output_dim; + + checkCuda(cudaMallocHost(&confidences_h, sizeof(float) * odim.tot())); + checkCuda(cudaMalloc(&tmpInputData_d, sizeof(float) * odim.tot())); + checkCuda(cudaMalloc(&tmpOutData_d, sizeof(float) * odim.w*odim.h)); + checkCuda(cudaMallocHost(&tmpOutData_h, sizeof(float) * odim.w*odim.h)); segmented.resize(nBatches); masks.resize(nBatches); @@ -208,7 +210,7 @@ class SegmentationNN { /** * Method to draw boundixg boxes and labels on a frame. */ - void draw(const int cur_batches=1) { + cv::Mat draw(const int cur_batches=1) { for(int i=0; i #include #include #include @@ -9,6 +10,8 @@ #include #include #include +#include + #include "tkdnn.h" @@ -36,4 +39,6 @@ void topKxyAddOffset(int * ids_begin, const int K, const int size, int *intxs_be void bboxes(int * ids_begin, const int K, const int size, float *xs_begin, float *ys_begin, dnnType *src_begin, float *bbx0, float *bbx1, float *bby0, float *bby1, float *src_out, int *ids_out); +void maxElem(dnnType *src_begin, dnnType *dst_begin, const int c, const int h, const int w); + #endif //KERNELSTHRUST_H \ No newline at end of file diff --git a/src/kernels/postprocessing.cu b/src/kernels/postprocessing.cu index 3510200..53234a2 100644 --- a/src/kernels/postprocessing.cu +++ b/src/kernels/postprocessing.cu @@ -34,6 +34,25 @@ void sortAndTopKonDevice(dnnType *src_begin, int *idsrc, float *topk_scores, int sortAndTopK_kernel<<>>(src_begin, idsrc, topk_scores, topk_inds, topk_ys, topk_xs, size, K); } +__global__ +void maxElem_kernel(float *src_begin, float *dst_begin, const int n_classes, const int size){ + int i = blockDim.x*blockIdx.x + threadIdx.x; + if (i > size) + return; + + thrust::device_ptr dPbeg ( &src_begin[i*n_classes] ) ; + thrust::device_ptr dPend = dPbeg + n_classes; + thrust::device_ptr result = thrust::max_element(thrust::device,dPbeg, dPend); + + dst_begin[i] = result - dPbeg; +} + +void maxElem(dnnType *src_begin, dnnType *dst_begin, const int c, const int h, const int w){ + int blocks = (h*w)/32+1; + int threads = 32; + maxElem_kernel<<>>(src_begin, dst_begin, c, h*w); +} + void topKxyclasses(int *ids_begin, int *ids_end, const int K, const int size, const int wh, int *clses, int *xs, int *ys){ thrust::transform(thrust::device, ids_begin, ids_end, thrust::make_constant_iterator(wh), clses, thrust::divides()); thrust::transform(thrust::device, ids_begin, ids_end, thrust::make_constant_iterator(wh), ids_begin, thrust::modulus()); From 3bf954750273e342fa7d42ab95fe10068f070fab Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Mon, 29 Jun 2020 15:11:52 +0200 Subject: [PATCH 06/17] Improve preprocessing Signed-off-by: Micaela Verucchi --- include/tkDNN/SegmentationNN.h | 46 ++++++++++++++++++---------------- 1 file changed, 24 insertions(+), 22 deletions(-) diff --git a/include/tkDNN/SegmentationNN.h b/include/tkDNN/SegmentationNN.h index 0289d61..87c7b8c 100644 --- a/include/tkDNN/SegmentationNN.h +++ b/include/tkDNN/SegmentationNN.h @@ -13,8 +13,6 @@ #include #include - - #include "tkdnn.h" #include "NetworkViz.h" #include "kernelsThrust.h" @@ -38,6 +36,8 @@ class SegmentationNN { float *tmpOutData_d; float *tmpOutData_h; + float *mean_d, *stddev_d; + cublasHandle_t cublasHandle; /** @@ -47,23 +47,10 @@ class SegmentationNN { * @param bi batch index */ void preprocess(cv::Mat &frame, const int bi=0) { - frame.convertTo(frame, CV_32FC3, 1 / 255.0, 0); - - cv::split(frame, bgr); - float mean[] = {0.485, 0.456, 0.406}; - float stddev[] = {0.229, 0.224, 0.225}; - for(int i=0; i<3; i++){ - bgr[2-i] -= mean[i]; - bgr[2-i] /= stddev[i]; - } - cv::merge(bgr, 3, frame); - - int crop_size = netRT->input_dim.w; int H = frame.rows; int W = frame.cols; cv::Mat frame_cropped; - cv::Mat mask(frame.size(), CV_8UC3, cv::Scalar(255,255,255)); if(H != W){ @@ -81,17 +68,22 @@ class SegmentationNN { } } - resize(frame_cropped, frame_cropped, cv::Size(netRT->input_dim.w, netRT->input_dim.h)); - resize(mask, mask, cv::Size(netRT->input_dim.w, netRT->input_dim.h)); - masks[bi] = mask.clone(); + tk::dnn::dataDim_t idim = netRT->input_dim; + + resize(frame_cropped, frame_cropped, cv::Size(idim.w, idim.h)); + resize(mask, mask, cv::Size(idim.w, idim.h)); + masks[bi] = mask; cv::split(frame_cropped, bgr); - for (int i = 0; i < netRT->input_dim.c; i++){ + for (int i = 0; i < idim.c; i++){ int idx = i * frame_cropped.rows * frame_cropped.cols; - int ch = netRT->input_dim.c-1 -i; - memcpy((void *)&input[idx + netRT->input_dim.tot()*bi], (void *)bgr[ch].data, frame_cropped.rows * frame_cropped.cols * sizeof(dnnType)); + int ch = idim.c-1 -i; + memcpy((void *)&input[idx + idim.tot()*bi], (void *)bgr[ch].data, frame_cropped.rows * frame_cropped.cols * sizeof(dnnType)); } - checkCuda(cudaMemcpyAsync(input_d+ netRT->input_dim.tot()*bi, input + netRT->input_dim.tot()*bi, netRT->input_dim.tot() * sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream)); + + checkCuda(cudaMemcpyAsync(input_d+ idim.tot()*bi, input + idim.tot()*bi, idim.tot() * sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream)); + + normalize(input_d + idim.tot()*bi, idim.c, idim.h, idim.w, mean_d, stddev_d); } /** @@ -157,6 +149,16 @@ class SegmentationNN { segmented.resize(nBatches); masks.resize(nBatches); + + std::vector mean = {0.485, 0.456, 0.406}; + std::vector stddev = {0.229, 0.224, 0.225}; + + checkCuda(cudaMalloc(&mean_d, sizeof(float) * mean.size())); + checkCuda(cudaMalloc(&stddev_d, sizeof(float) * stddev.size())); + + checkCuda(cudaMemcpyAsync(mean_d, mean.data(), mean.size() * sizeof(float), cudaMemcpyHostToDevice, netRT->stream)); + checkCuda(cudaMemcpyAsync(stddev_d, stddev.data(), stddev.size() * sizeof(float), cudaMemcpyHostToDevice, netRT->stream)); + } /** From a5cc4e3edada046902cb9fd3eea0ad1efe59970a Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Wed, 1 Jul 2020 10:53:24 +0200 Subject: [PATCH 07/17] Add berkeley test, add weights for shelfnets Signed-off-by: Micaela Verucchi --- CMakeLists.txt | 3 + demo/demo/seg_demo.cpp | 9 +- tests/shelfnet/shelfnet.cpp | 2 +- tests/shelfnet/shelfnet_berkeley.cpp | 295 +++++++++++++++++++++++++++ 4 files changed, 305 insertions(+), 4 deletions(-) create mode 100644 tests/shelfnet/shelfnet_berkeley.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 88e94e3..9e9c27c 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -107,6 +107,9 @@ target_link_libraries(test_dla34_cnet tkDNN) add_executable(test_shelfnet tests/shelfnet/shelfnet.cpp) target_link_libraries(test_shelfnet tkDNN) +add_executable(test_shelfnet_berkeley tests/shelfnet/shelfnet_berkeley.cpp) +target_link_libraries(test_shelfnet_berkeley tkDNN) + # DEMOS add_executable(test_rtinference tests/test_rtinference/rtinference.cpp) target_link_libraries(test_rtinference tkDNN) diff --git a/demo/demo/seg_demo.cpp b/demo/demo/seg_demo.cpp index c1ee43b..0737ad7 100644 --- a/demo/demo/seg_demo.cpp +++ b/demo/demo/seg_demo.cpp @@ -29,9 +29,12 @@ int main(int argc, char *argv[]) { int n_batch = 1; if(argc > 3) n_batch = atoi(argv[3]); - bool show = false; + int n_classes = 19; if(argc > 4) - show = atoi(argv[4]); + n_classes = atoi(argv[4]); + bool show = false; + if(argc > 5) + show = atoi(argv[5]); if(n_batch < 1 || n_batch > 64) FatalError("Batch dim not supported"); @@ -39,7 +42,7 @@ int main(int argc, char *argv[]) { if(!show) SAVE_RESULT = true; - int n_classes = 19; + tk::dnn::SegmentationNN segNN; segNN.init(net, n_classes, n_batch); diff --git a/tests/shelfnet/shelfnet.cpp b/tests/shelfnet/shelfnet.cpp index 7cc1b42..48cad04 100644 --- a/tests/shelfnet/shelfnet.cpp +++ b/tests/shelfnet/shelfnet.cpp @@ -83,7 +83,7 @@ const char *trans[] = { int main() { - // downloadWeightsifDoNotExist(input_bin, "shelfnet", "https://cloud.hipert.unimore.it/s/x4ZfxBKN23zAJQp/download"); + downloadWeightsifDoNotExist(input_bin, "shelfnet", "https://cloud.hipert.unimore.it/s/mEDZMRJaGCFWSJF/download"); int classes = 19; diff --git a/tests/shelfnet/shelfnet_berkeley.cpp b/tests/shelfnet/shelfnet_berkeley.cpp new file mode 100644 index 0000000..fe81191 --- /dev/null +++ b/tests/shelfnet/shelfnet_berkeley.cpp @@ -0,0 +1,295 @@ +#include +#include +#include + +#include "tkdnn.h" +#include "NetworkViz.h" + + +const char *input_bin = "shelfnet_berkeley/debug/input.bin"; + +const char *backbone[] = { + "shelfnet_berkeley/layers/backbone-conv1.bin", + "shelfnet_berkeley/layers/backbone-layer1-0-conv1.bin", + "shelfnet_berkeley/layers/backbone-layer1-0-conv2.bin", + "shelfnet_berkeley/layers/backbone-layer1-1-conv1.bin", + "shelfnet_berkeley/layers/backbone-layer1-1-conv2.bin", + "shelfnet_berkeley/layers/backbone-layer2-0-conv1.bin", + "shelfnet_berkeley/layers/backbone-layer2-0-conv2.bin", + "shelfnet_berkeley/layers/backbone-layer2-0-downsample-0.bin", + "shelfnet_berkeley/layers/backbone-layer2-1-conv1.bin", + "shelfnet_berkeley/layers/backbone-layer2-1-conv2.bin", + "shelfnet_berkeley/layers/backbone-layer3-0-conv1.bin", + "shelfnet_berkeley/layers/backbone-layer3-0-conv2.bin", + "shelfnet_berkeley/layers/backbone-layer3-0-downsample-0.bin", + "shelfnet_berkeley/layers/backbone-layer3-1-conv1.bin", + "shelfnet_berkeley/layers/backbone-layer3-1-conv2.bin", + "shelfnet_berkeley/layers/backbone-layer4-0-conv1.bin", + "shelfnet_berkeley/layers/backbone-layer4-0-conv2.bin", + "shelfnet_berkeley/layers/backbone-layer4-0-downsample-0.bin", + "shelfnet_berkeley/layers/backbone-layer4-1-conv1.bin", + "shelfnet_berkeley/layers/backbone-layer4-1-conv2.bin"}; + +const char *conv_out[] = { + "shelfnet_berkeley/layers/conv_out-conv-conv.bin", + "shelfnet_berkeley/layers/conv_out-conv_out.bin", + "shelfnet_berkeley/layers/conv_out16-conv-conv.bin", + "shelfnet_berkeley/layers/conv_out16-conv_out.bin", + "shelfnet_berkeley/layers/conv_out32-conv-conv.bin", + "shelfnet_berkeley/layers/conv_out32-conv_out.bin" + }; + +const char *decoder[] = { + "shelfnet_berkeley/layers/decoder-bottom-conv1.bin", + "shelfnet_berkeley/layers/decoder-bottom-conv12.bin", + "shelfnet_berkeley/layers/decoder-up_conv_list-0-conv-conv.bin", + "shelfnet_berkeley/layers/decoder-up_conv_list-0-conv_atten.bin", + "shelfnet_berkeley/layers/decoder-up_dense_list-0-conv.bin", + "shelfnet_berkeley/layers/decoder-up_conv_list-1-conv-conv.bin", + "shelfnet_berkeley/layers/decoder-up_conv_list-1-conv_atten.bin", + "shelfnet_berkeley/layers/decoder-up_dense_list-1-conv.bin" + }; + + +const char *ladder[] = { + "shelfnet_berkeley/layers/ladder-inconv-conv1.bin", + "shelfnet_berkeley/layers/ladder-inconv-conv12.bin", + "shelfnet_berkeley/layers/ladder-down_module_list-0-conv1.bin", + "shelfnet_berkeley/layers/ladder-down_module_list-0-conv12.bin", + "shelfnet_berkeley/layers/ladder-down_conv_list-0.bin", + + "shelfnet_berkeley/layers/ladder-down_module_list-1-conv1.bin", + "shelfnet_berkeley/layers/ladder-down_module_list-1-conv12.bin", + "shelfnet_berkeley/layers/ladder-down_conv_list-1.bin", + + "shelfnet_berkeley/layers/ladder-bottom-conv1.bin", + "shelfnet_berkeley/layers/ladder-bottom-conv12.bin", + + + + "shelfnet_berkeley/layers/ladder-up_conv_list-0-conv-conv.bin", + "shelfnet_berkeley/layers/ladder-up_conv_list-0-conv_atten.bin", + "shelfnet_berkeley/layers/ladder-up_dense_list-0-conv.bin", + + + "shelfnet_berkeley/layers/ladder-up_conv_list-1-conv-conv.bin", + "shelfnet_berkeley/layers/ladder-up_conv_list-1-conv_atten.bin", + "shelfnet_berkeley/layers/ladder-up_dense_list-1-conv.bin"}; + +const char *trans[] = { + "shelfnet_berkeley/layers/trans1-conv.bin", + "shelfnet_berkeley/layers/trans2-conv.bin", + "shelfnet_berkeley/layers/trans3-conv.bin"}; +int main() +{ + + downloadWeightsifDoNotExist(input_bin, "shelfnet_berkeley", "https://cloud.hipert.unimore.it/s/m92e7QdD9gYMF7f/download"); + + int classes = 20; + + // Network layout + tk::dnn::dataDim_t dim(1, 3, 1024, 1024, 1); + tk::dnn::Network net(dim); + + int bi = 0, di = 0, li = 0, ci = 0; + new tk::dnn::Conv2d(&net, 64, 7, 7, 2, 2, 3, 3, backbone[bi++], true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + tk::dnn::Layer* last = new tk::dnn::Pooling (&net, 3, 3, 2, 2, 1, 1, tk::dnn::POOLING_MAX); + + + + for(int i=0; i<2; ++i){ + new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, backbone[bi++], true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, backbone[bi++], true); + new tk::dnn::Shortcut(&net, last); + last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); + } + + std::vector features; + for(int i=0;i<3;++i){ + int out_channel = pow(2,7+i); + std::cout< up_out; + //bottom + new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, decoder[di++], true, false, 1, true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, decoder[di++], true, false, 1, true); + new tk::dnn::Shortcut(&net, last); + last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); + up_out.push_back(last); + + for(int i=0; i<2; ++i){ + int out_channel = pow(2,7-i); + //up-conv + std::cout<output_dim.w, last->output_dim.h, last->output_dim.w, last->output_dim.h, 0, 0, tk::dnn::POOLING_AVERAGE); + new tk::dnn::Conv2d (&net, out_channel, 1, 1, 1, 1, 0, 0, decoder[di++], true); + + tk::dnn::Layer* act = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_SIGMOID); + new tk::dnn::Route(&net, &last, 1); + new tk::dnn::Shortcut(&net, act, true); + + //interpolate + new tk::dnn::Resize(&net, 1,2,2); + new tk::dnn::Shortcut(&net, features[1-i]); + + //up-dense + new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, decoder[di++], true); + last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + up_out.push_back(last); + } + + //LADDER + + std::vector down_out; + new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); + new tk::dnn::Shortcut(&net, last); + new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); + + for(int i=0; i<2;++i){ + int out_channel = pow(2,6+i); + tk::dnn::Layer* l_last = new tk::dnn::Shortcut(&net, up_out[2-i]); + + new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); + new tk::dnn::Shortcut(&net, l_last); + l_last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); + down_out.push_back(l_last); + + new tk::dnn::Conv2d (&net, out_channel*2, 3, 3, 2, 2, 1, 1, ladder[li++], false); + last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.0f); //should be ReLU + } + + new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); + new tk::dnn::Shortcut(&net, last); + last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); + up_out.clear(); + up_out.push_back(last); + + for(int i=0; i<2; ++i){ + int out_channel = pow(2,7-i); + //up-conv + new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true); + last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + + new tk::dnn::Pooling(&net, last->output_dim.w, last->output_dim.h, last->output_dim.w, last->output_dim.h, 0, 0, tk::dnn::POOLING_AVERAGE); + new tk::dnn::Conv2d (&net, out_channel, 1, 1, 1, 1, 0, 0, ladder[li++], true); + + tk::dnn::Layer* act = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_SIGMOID); + new tk::dnn::Route(&net, &last, 1); + new tk::dnn::Shortcut(&net, act, true); + + //interpolate + new tk::dnn::Resize(&net, 1,2,2); + new tk::dnn::Shortcut(&net, down_out[1-i]); + + // //up-dense + new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true); + last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + up_out.push_back(last); + } + + + // for(int i=2;i>=0;--i){ + // new tk::dnn::Route(&net, &up_out[i], 1); + new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, conv_out[ci++], true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + new tk::dnn::Conv2d (&net, classes, 3, 3, 1, 1, 1, 1, conv_out[ci++], false); + /*up_out[i] =*/ new tk::dnn::Resize(&net, classes, net.input_dim.h, net.input_dim.w, true, tk::dnn::ResizeMode_t::LINEAR); + // } + + new tk::dnn::Softmax(&net); + + const char *output_bin = "shelfnet_berkeley/debug/softmax.bin"; + + // Load input + dnnType *data; + dnnType *input_h; + readBinaryFile(input_bin, dim.tot(), &input_h, &data); + std::cout<<"Input:"< batch_frame; - std::vector batch_dnn_input; - int height = 0, width = 0; - - while(gRun) { - batch_dnn_input.clear(); - batch_frame.clear(); - - for(int bi=0; bi< n_batch; ++bi){ + cv::Mat frame; + while(gRun) { cap >> frame; if(!frame.data) break; height = frame.rows; width = frame.cols; - batch_frame.push_back(frame); - // this will be resized to the net format - batch_dnn_input.push_back(frame.clone()); - } - if(!frame.data) - break; - - //inference - segNN.update(batch_dnn_input, n_batch); - frame = segNN.draw(); + //inference + segNN.updateOriginal(frame); + if(show) + segNN.draw(); - if(n_batch == 1 && SAVE_RESULT) - resultVideo << frame; + if(SAVE_RESULT) + resultVideo << segNN.segmented[0]; + } } std::cout<<"segmentation end\n"; double mean = 0, mean_pre = 0, mean_post = 0; std::cout< splitted_frames; + int H, W, net_H, net_W; + int top = 0, bottom = 0, left = 0, right = 0; + std::vector> pos; + + { + TKDNN_TSTART + cv::Size original_size = frame.size(); + + frame.convertTo(frame, CV_32FC3, 1 / 255.0, 0); + H = frame.rows; + W = frame.cols; + net_H = netRT->input_dim.h; + net_W = netRT->input_dim.w; + + cv::Mat frame_cropped; + + if( H <= net_H && W <= net_W ){ // smaller size wrt network + top = (net_H - H)/2; + bottom = net_H - H - top ; + left = (net_W - W)/2; + right = net_W - W - left ; + cv::copyMakeBorder(frame, frame_cropped, top, bottom, left, right, cv::BORDER_CONSTANT, cv::Scalar(0,0,0) ); + splitted_frames.push_back(frame_cropped); + } + else{ //bigger size wrt network + + + if(H < net_H || W < net_W){ + if(H < net_H){ + top = (net_H - H)/2; + bottom = net_H - H - top ; + } + else{ + left = (net_W - W)/2; + right = net_W - W - left ; + } + cv::copyMakeBorder(frame, frame_cropped, top, bottom, left, right, cv::BORDER_CONSTANT, cv::Scalar(0,0,0)); + } + + for(int x=0; x+net_W<=W ;){ + for(int y=0; y+net_H <=H ; ){ + cv::Rect roi(x, y, net_W, net_H); + cv::Mat image_roi = frame(roi); + splitted_frames.push_back(image_roi); + pos.push_back(std::make_pair(x,y)); + + y += net_H; + if(y == H) + break; + if(y + net_H > H) y = H - net_H; + } + x += net_W; + if(x == W) + break; + if(x + net_W > W) x = W - net_W; + } + } + + tk::dnn::dataDim_t idim = netRT->input_dim; + + if(splitted_frames.size()> nBatches) + FatalError(std::to_string(splitted_frames.size()) + " min batches required"); + + for(int bi=0; bistream)); + normalize(input_d + idim.tot()*bi, idim.c, idim.h, idim.w, mean_d, stddev_d); + } + TKDNN_TSTOP + stats_pre.push_back(t_ns); + } + + tk::dnn::dataDim_t dim = netRT->input_dim; + dim.n = splitted_frames.size(); + { + if(TKDNN_VERBOSE) dim.print(); + TKDNN_TSTART + netRT->infer(dim, input_d); + TKDNN_TSTOP + if(TKDNN_VERBOSE) dim.print(); + stats.push_back(t_ns); + } + + dataDim_t odim = netRT->output_dim; + + std::vector out_img; + + { + TKDNN_TSTART + + for(int bi=0; bibuffersRT[1]+ netRT->buffersDIM[1].tot()*bi; + + matrixTranspose(cublasHandle, rt_out, tmpInputData_d, odim.c, odim.w*odim.h); + maxElem(tmpInputData_d, tmpOutData_d, odim.c, odim.h, odim.w); + checkCuda(cudaMemcpy(tmpOutData_h, tmpOutData_d, odim.w*odim.h * sizeof(float), cudaMemcpyDeviceToHost)); + + dataDim_t vdim = odim; + vdim.c = 1; + + cv::Mat colored; + + if(apply_colormap) + colored = vizData2Mat(tmpOutData_h, vdim, 1024, 0, 18); + else{ + cv::Mat colored_fp32 (cv::Size(odim.w, odim.h),CV_32FC1, tmpOutData_h); + colored_fp32.convertTo(colored, CV_8UC1); + } + out_img.push_back(colored); + } + + + cv::Mat seg(frame.size(), out_img[0].type()); + if(out_img.size() == 1) + { + cv::Rect roi(left, top, W, H); + seg = out_img[0](roi); + } + else{ + int bi=0; + + if(top == 0 && left == 0){ + + for(int i=0; i Date: Tue, 24 Nov 2020 12:37:33 +0100 Subject: [PATCH 11/17] Add shelfnet_mapillary, README_seg, resize of input Signed-off-by: Micaela Verucchi --- CMakeLists.txt | 3 + README_seg.md | 56 +++++ demo/demo/seg_demo.cpp | 48 ++++- include/tkDNN/NetworkViz.h | 4 +- include/tkDNN/SegmentationNN.h | 8 +- src/NetworkViz.cpp | 108 +++++++++- tests/shelfnet/shelfnet_mapillary.cpp | 295 ++++++++++++++++++++++++++ 7 files changed, 498 insertions(+), 24 deletions(-) create mode 100644 README_seg.md create mode 100644 tests/shelfnet/shelfnet_mapillary.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 9e9c27c..c292300 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -110,6 +110,9 @@ target_link_libraries(test_shelfnet tkDNN) add_executable(test_shelfnet_berkeley tests/shelfnet/shelfnet_berkeley.cpp) target_link_libraries(test_shelfnet_berkeley tkDNN) +add_executable(test_shelfnet_mapillary tests/shelfnet/shelfnet_mapillary.cpp) +target_link_libraries(test_shelfnet_mapillary tkDNN) + # DEMOS add_executable(test_rtinference tests/test_rtinference/rtinference.cpp) target_link_libraries(test_rtinference tkDNN) diff --git a/README_seg.md b/README_seg.md new file mode 100644 index 0000000..a9afcd0 --- /dev/null +++ b/README_seg.md @@ -0,0 +1,56 @@ +# Semantic Segmentation with tkDNN + +Currently tkDNN supports only ShelfNet as semantic segmentation network. + +## Export weights from Shelfnet +To get the weights needed to run Mobilenet tests use [this](https://git.hipert.unimore.it/mverucchi/shelfnet) fork of a Pytorch implementation of Shelfnet network. + +``` +git clone https://git.hipert.unimore.it/mverucchi/shelfnet +cd shelfnet +cd ShelfNet18_realtime +conda env create --file shelfnet_env.yml +conda activate shelfnet +mkdir layer debug +python export.py +``` + + +## Run the demo + +To run the semantic segmentation demo follow these steps (example with shelfnet_mapillary): +``` +rm shelfnet_mapillary_fp32.rt # be sure to delete(or move) old tensorRT files +export TKDNN_BATCHSIZE=4 # be sure you have batch size > than 1 if you want to run inference on images bigger than 1024 +./test_shelfnet_mapillary # run the yolo test (is slow) +./demo shelfnet_mapillary_fp32.rt ../demo/yolo_test.mp4 1 15 +``` +In general the demo program takes the following parameters: +``` +./seg_demo +``` +where +* `````` is the rt file generated by a test +* ```<``` is the path to a video file or a camera input +* `````` number of batches to use in inference (N.B. you should first export TKDNN_BATCHSIZE to the required n_batches and create again the rt file for the network). +* ``````is the number of classes the network is trained on +* `````` if set to 0 the demo will not resize the input frames, but use it as it is, otherwise it will resize it. +* `````` is `````` is set to 1, then the input frames will be proportionally resized using `````` as width baseline. +* `````` if set to 0 the demo will not show the visualization but save the video into result.mp4 (if n-batches ==1) +* `````` if set to 0 (deafult) the demo will run, otherwise the evaluation of a dataset will run and the output of the segmentation will be saved. Attention: this is under development and paths are embedded, so change them in the code in advance. + +N.b. By default it is used FP32 inference + + + +## Existing tests and supported networks + +| Test Name | Network | Dataset | N Classes | Input size | Weights | +| :---------------- | :-------------------------------------------- | :-----------------------------------------------------------: | :-------: | :-----------: | :------------------------------------------------------------------------ | +| shelfnet | ShelfNet18_realtime1 | [Cityscapes](https://www.cityscapes-dataset.com/) | 19 | 1024x1024 | [weights](https://cloud.hipert.unimore.it/s/mEDZMRJaGCFWSJF/download) | +| shelfnet_berkeley | ShelfNet18_realtime1 | [DeepDrive](https://bdd-data.berkeley.edu/) | 20 | 1024x1024 | [weights](https://cloud.hipert.unimore.it/s/m92e7QdD9gYMF7f/download) | +| shelfnet_mapillary | ShelfNet18_realtime1 | [Mapillary Vistas](https://www.mapillary.com/dataset/vistas?pKey=aFWuj_m4nGoq3-tDz5KAqQ)* | 15 | 1024x1024 | [weights](https://cloud.hipert.unimore.it/s/6WnZCKLjik7xrny/download) | + +1. Zhuang, Juntang, et al. "ShelfNet for fast semantic segmentation." Proceedings of the IEEE International Conference on Computer Vision Workshops. 2019. + +*. Mapillary Vistas has originally 66 classes, but we reduced them to 15 to improve the results on the categories of our interest. \ No newline at end of file diff --git a/demo/demo/seg_demo.cpp b/demo/demo/seg_demo.cpp index 35b4fa8..bd31a1a 100644 --- a/demo/demo/seg_demo.cpp +++ b/demo/demo/seg_demo.cpp @@ -48,22 +48,37 @@ int main(int argc, char *argv[]) { int n_classes = 19; if(argc > 4) n_classes = atoi(argv[4]); - bool show = true; + bool resize = false; if(argc > 5) - show = atoi(argv[5]); - bool write_pred = false; + resize = atoi(argv[5]); + int baseline_resize = 1024; if(argc > 6) - write_pred = atoi(argv[6]); - + baseline_resize = atoi(argv[6]); + bool show = true; + if(argc > 7) + show = atoi(argv[7]); + bool write_pred = false; + if(argc > 8) + write_pred = atoi(argv[8]); + if(resize && (baseline_resize < 0 || baseline_resize > 5000)) + FatalError("Problem with baseline resize") if(n_batch < 1 || n_batch > 64) FatalError("Batch dim not supported"); + std::string net_name; + removePathAndExtension(net, net_name); + bool mapillary_15 = false; //TODO change me pls + if(n_classes == 15 && net_name == "shelfnet_mapillary_fp32") + mapillary_15 = true; + + //net initialization tk::dnn::SegmentationNN segNN; segNN.init(net, n_classes, n_batch); int height = 0, width = 0; - + int basewidth=baseline_resize, hsize; + if(write_pred){ std::string gt_folder = "../demo/CityScapes_val/images/"; std::string images_names = "../demo/CityScapes_val/all_images.txt"; @@ -85,9 +100,16 @@ int main(int argc, char *argv[]) { cv::VideoWriter resultVideo; if(SAVE_RESULT) { - int w = cap.get(cv::CAP_PROP_FRAME_WIDTH); - int h = cap.get(cv::CAP_PROP_FRAME_HEIGHT); - resultVideo.open("result.mp4", cv::VideoWriter::fourcc('M','P','4','V'), 30, cv::Size(1024, 1024)); + int w,h; + if(resize){ + w = basewidth; + h = int((float(cap.get(cv::CAP_PROP_FRAME_HEIGHT))*float(basewidth/float(cap.get(cv::CAP_PROP_FRAME_WIDTH))))); + } + else{ + w = cap.get(cv::CAP_PROP_FRAME_WIDTH); + h = cap.get(cv::CAP_PROP_FRAME_HEIGHT); + } + resultVideo.open("result.mp4", cv::VideoWriter::fourcc('M','P','4','V'), 30, cv::Size(w, h)); } cv::Mat frame; @@ -95,11 +117,17 @@ int main(int argc, char *argv[]) { cap >> frame; if(!frame.data) break; + + if(resize){ + hsize = int((float(frame.rows)*float(basewidth/float(frame.cols)))); + cv::resize(frame, frame, cv::Size(basewidth, hsize)); + } + height = frame.rows; width = frame.cols; //inference - segNN.updateOriginal(frame); + segNN.updateOriginal(frame, true, mapillary_15); if(show) segNN.draw(); diff --git a/include/tkDNN/NetworkViz.h b/include/tkDNN/NetworkViz.h index 94468fc..2e2c3f4 100644 --- a/include/tkDNN/NetworkViz.h +++ b/include/tkDNN/NetworkViz.h @@ -5,8 +5,8 @@ namespace tk { namespace dnn { -cv::Mat vizFloat2colorMap(cv::Mat map, double min=0, double max=0); -cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim, double min=0, double max=0); +cv::Mat vizFloat2colorMap(cv::Mat map, double min=0, double max=0, bool mapillary_15=false); +cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim, double min=0, double max=0, bool mapillary_15=false); cv::Mat vizLayer2Mat(tk::dnn::Network *net, int layer, int imgdim = 1000); }} diff --git a/include/tkDNN/SegmentationNN.h b/include/tkDNN/SegmentationNN.h index edede43..ec6ac79 100644 --- a/include/tkDNN/SegmentationNN.h +++ b/include/tkDNN/SegmentationNN.h @@ -97,7 +97,7 @@ class SegmentationNN { * * @param bi batch index */ - void postprocess(const int bi=0, bool appy_colormap = true) { + void postprocess(const int bi=0, bool appy_colormap = true, bool mapillary_15=false) { dnnType *rt_out = (dnnType *)netRT->buffersRT[1]+ netRT->buffersDIM[1].tot()*bi; dataDim_t odim = netRT->output_dim; @@ -112,7 +112,7 @@ class SegmentationNN { cv::Mat colored; if(appy_colormap) - colored = vizData2Mat(tmpOutData_h, vdim, 1024, 0, 18); + colored = vizData2Mat(tmpOutData_h, vdim, 1024, 0, classes, mapillary_15); else{ cv::Mat colored_fp32 (cv::Size(odim.w, odim.h),CV_32FC1, tmpOutData_h); colored_fp32.convertTo(colored, CV_8UC1); @@ -234,7 +234,7 @@ class SegmentationNN { } } - void updateOriginal(cv::Mat frame, bool apply_colormap=true){ + void updateOriginal(cv::Mat frame, bool apply_colormap=true, bool mapillary_15=false){ std::vector splitted_frames; int H, W, net_H, net_W; @@ -347,7 +347,7 @@ class SegmentationNN { cv::Mat colored; if(apply_colormap) - colored = vizData2Mat(tmpOutData_h, vdim, 1024, 0, 18); + colored = vizData2Mat(tmpOutData_h, vdim, 1024, 0, classes, mapillary_15); else{ cv::Mat colored_fp32 (cv::Size(odim.w, odim.h),CV_32FC1, tmpOutData_h); colored_fp32.convertTo(colored, CV_8UC1); diff --git a/src/NetworkViz.cpp b/src/NetworkViz.cpp index 842a26e..4a3f113 100644 --- a/src/NetworkViz.cpp +++ b/src/NetworkViz.cpp @@ -6,22 +6,114 @@ namespace tk { namespace dnn { -cv::Mat vizFloat2colorMap(cv::Mat map,double min, double max) { +cv::Mat mapillary_15_map(cv::Mat adjMap){ + + // cv::imshow("test", adjMap); + // cv::waitKey(0); + cv::Mat M1(1, 256, CV_8UC1), M2(1, 256, CV_8UC1), M3(1, 256, CV_8UC1); + + M3.at(0)=165; + M2.at(0)=42; + M1.at(0)=45; + + M3.at(1)=196; + M2.at(1)=196; + M1.at(1)=196; + + M3.at(2)=90; + M2.at(2)=120; + M1.at(2)=150; + + M3.at(3)=128; + M2.at(3)=64; + M1.at(3)=128; + + M3.at(4)=70; + M2.at(4)=70; + M1.at(4)=70; + + M3.at(5)=220; + M2.at(5)=20; + M1.at(5)=60; + + M3.at(6)=255; + M2.at(6)=255; + M1.at(6)=255; + + M3.at(7)=107; + M2.at(7)=142; + M1.at(7)=35; + + M3.at(8)=70; + M2.at(8)=130; + M1.at(8)=180; + + M3.at(9)=220; + M2.at(9)=220; + M1.at(9)=220; + + M3.at(10)=153; + M2.at(10)=153; + M1.at(10)=153; + + M3.at(11)=128; + M2.at(11)=128; + M1.at(11)=128; + + M3.at(12)=119; + M2.at(12)=11; + M1.at(12)=32; + + M3.at(13)=0; + M2.at(13)=0; + M1.at(13)=142; + + for(int i=14;i<256;i++) + { + M1.at(i)=0; + M2.at(i)=0; + M3.at(i)=0; + } + + cv::Mat r1,r2,r3; + + cv::LUT(adjMap,M1,r1); + cv::LUT(adjMap,M2,r2); + cv::LUT(adjMap,M3,r3); + + std::vector planes; + planes.push_back(r1); + planes.push_back(r2); + planes.push_back(r3); + + cv::Mat dst; + cv::merge(planes,dst); + return dst; + + +} + +cv::Mat vizFloat2colorMap(cv::Mat map,double min, double max, bool mapillary_15) { if(min == 0 && max == 0) cv::minMaxIdx(map, &min, &max); cv::Mat adjMap; - // expand your range to 0..255. Similar to histEq(); - map.convertTo(adjMap,CV_8UC1, 255 / (max-min), -min); - //return adjMap; - cv::Mat falseColorsMap; - applyColorMap(adjMap, falseColorsMap, cv::COLORMAP_VIRIDIS); + + if(mapillary_15){ + map.convertTo(adjMap,CV_8UC1); + falseColorsMap = mapillary_15_map(adjMap); + } + else{ + // expand your range to 0..255. Similar to histEq(); + map.convertTo(adjMap,CV_8UC1, 255 / (max-min), -min); + applyColorMap(adjMap, falseColorsMap, cv::COLORMAP_JET); + } return falseColorsMap; } -cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim, double min, double max) { +cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim, double min, double max, bool mapillary_15) { dnnType *data = nullptr; // copy to CPU @@ -37,7 +129,7 @@ cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim, doub cv::Mat grid = cv::Mat(gridSize, CV_8UC3, cv::Scalar(0)); for(int i=0; i +#include +#include + +#include "tkdnn.h" +#include "NetworkViz.h" + + +const char *input_bin = "shelfnet_mapillary/debug/input.bin"; + +const char *backbone[] = { + "shelfnet_mapillary/layers/backbone-conv1.bin", + "shelfnet_mapillary/layers/backbone-layer1-0-conv1.bin", + "shelfnet_mapillary/layers/backbone-layer1-0-conv2.bin", + "shelfnet_mapillary/layers/backbone-layer1-1-conv1.bin", + "shelfnet_mapillary/layers/backbone-layer1-1-conv2.bin", + "shelfnet_mapillary/layers/backbone-layer2-0-conv1.bin", + "shelfnet_mapillary/layers/backbone-layer2-0-conv2.bin", + "shelfnet_mapillary/layers/backbone-layer2-0-downsample-0.bin", + "shelfnet_mapillary/layers/backbone-layer2-1-conv1.bin", + "shelfnet_mapillary/layers/backbone-layer2-1-conv2.bin", + "shelfnet_mapillary/layers/backbone-layer3-0-conv1.bin", + "shelfnet_mapillary/layers/backbone-layer3-0-conv2.bin", + "shelfnet_mapillary/layers/backbone-layer3-0-downsample-0.bin", + "shelfnet_mapillary/layers/backbone-layer3-1-conv1.bin", + "shelfnet_mapillary/layers/backbone-layer3-1-conv2.bin", + "shelfnet_mapillary/layers/backbone-layer4-0-conv1.bin", + "shelfnet_mapillary/layers/backbone-layer4-0-conv2.bin", + "shelfnet_mapillary/layers/backbone-layer4-0-downsample-0.bin", + "shelfnet_mapillary/layers/backbone-layer4-1-conv1.bin", + "shelfnet_mapillary/layers/backbone-layer4-1-conv2.bin"}; + +const char *conv_out[] = { + "shelfnet_mapillary/layers/conv_out-conv-conv.bin", + "shelfnet_mapillary/layers/conv_out-conv_out.bin", + "shelfnet_mapillary/layers/conv_out16-conv-conv.bin", + "shelfnet_mapillary/layers/conv_out16-conv_out.bin", + "shelfnet_mapillary/layers/conv_out32-conv-conv.bin", + "shelfnet_mapillary/layers/conv_out32-conv_out.bin" + }; + +const char *decoder[] = { + "shelfnet_mapillary/layers/decoder-bottom-conv1.bin", + "shelfnet_mapillary/layers/decoder-bottom-conv12.bin", + "shelfnet_mapillary/layers/decoder-up_conv_list-0-conv-conv.bin", + "shelfnet_mapillary/layers/decoder-up_conv_list-0-conv_atten.bin", + "shelfnet_mapillary/layers/decoder-up_dense_list-0-conv.bin", + "shelfnet_mapillary/layers/decoder-up_conv_list-1-conv-conv.bin", + "shelfnet_mapillary/layers/decoder-up_conv_list-1-conv_atten.bin", + "shelfnet_mapillary/layers/decoder-up_dense_list-1-conv.bin" + }; + + +const char *ladder[] = { + "shelfnet_mapillary/layers/ladder-inconv-conv1.bin", + "shelfnet_mapillary/layers/ladder-inconv-conv12.bin", + "shelfnet_mapillary/layers/ladder-down_module_list-0-conv1.bin", + "shelfnet_mapillary/layers/ladder-down_module_list-0-conv12.bin", + "shelfnet_mapillary/layers/ladder-down_conv_list-0.bin", + + "shelfnet_mapillary/layers/ladder-down_module_list-1-conv1.bin", + "shelfnet_mapillary/layers/ladder-down_module_list-1-conv12.bin", + "shelfnet_mapillary/layers/ladder-down_conv_list-1.bin", + + "shelfnet_mapillary/layers/ladder-bottom-conv1.bin", + "shelfnet_mapillary/layers/ladder-bottom-conv12.bin", + + + + "shelfnet_mapillary/layers/ladder-up_conv_list-0-conv-conv.bin", + "shelfnet_mapillary/layers/ladder-up_conv_list-0-conv_atten.bin", + "shelfnet_mapillary/layers/ladder-up_dense_list-0-conv.bin", + + + "shelfnet_mapillary/layers/ladder-up_conv_list-1-conv-conv.bin", + "shelfnet_mapillary/layers/ladder-up_conv_list-1-conv_atten.bin", + "shelfnet_mapillary/layers/ladder-up_dense_list-1-conv.bin"}; + +const char *trans[] = { + "shelfnet_mapillary/layers/trans1-conv.bin", + "shelfnet_mapillary/layers/trans2-conv.bin", + "shelfnet_mapillary/layers/trans3-conv.bin"}; +int main() +{ + + downloadWeightsifDoNotExist(input_bin, "shelfnet_mapillary", "https://cloud.hipert.unimore.it/s/6WnZCKLjik7xrny/download"); + + int classes = 15; + + // Network layout + tk::dnn::dataDim_t dim(1, 3, 1024, 1024, 1); + tk::dnn::Network net(dim); + + int bi = 0, di = 0, li = 0, ci = 0; + new tk::dnn::Conv2d(&net, 64, 7, 7, 2, 2, 3, 3, backbone[bi++], true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + tk::dnn::Layer* last = new tk::dnn::Pooling (&net, 3, 3, 2, 2, 1, 1, tk::dnn::POOLING_MAX); + + + + for(int i=0; i<2; ++i){ + new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, backbone[bi++], true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, backbone[bi++], true); + new tk::dnn::Shortcut(&net, last); + last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); + } + + std::vector features; + for(int i=0;i<3;++i){ + int out_channel = pow(2,7+i); + std::cout< up_out; + //bottom + new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, decoder[di++], true, false, 1, true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, decoder[di++], true, false, 1, true); + new tk::dnn::Shortcut(&net, last); + last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); + up_out.push_back(last); + + for(int i=0; i<2; ++i){ + int out_channel = pow(2,7-i); + //up-conv + std::cout<output_dim.w, last->output_dim.h, last->output_dim.w, last->output_dim.h, 0, 0, tk::dnn::POOLING_AVERAGE); + new tk::dnn::Conv2d (&net, out_channel, 1, 1, 1, 1, 0, 0, decoder[di++], true); + + tk::dnn::Layer* act = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_SIGMOID); + new tk::dnn::Route(&net, &last, 1); + new tk::dnn::Shortcut(&net, act, true); + + //interpolate + new tk::dnn::Resize(&net, 1,2,2); + new tk::dnn::Shortcut(&net, features[1-i]); + + //up-dense + new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, decoder[di++], true); + last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + up_out.push_back(last); + } + + //LADDER + + std::vector down_out; + new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); + new tk::dnn::Shortcut(&net, last); + new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); + + for(int i=0; i<2;++i){ + int out_channel = pow(2,6+i); + tk::dnn::Layer* l_last = new tk::dnn::Shortcut(&net, up_out[2-i]); + + new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); + new tk::dnn::Shortcut(&net, l_last); + l_last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); + down_out.push_back(l_last); + + new tk::dnn::Conv2d (&net, out_channel*2, 3, 3, 2, 2, 1, 1, ladder[li++], false); + last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.0f); //should be ReLU + } + + new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true); + new tk::dnn::Shortcut(&net, last); + last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); + up_out.clear(); + up_out.push_back(last); + + for(int i=0; i<2; ++i){ + int out_channel = pow(2,7-i); + //up-conv + new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true); + last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + + new tk::dnn::Pooling(&net, last->output_dim.w, last->output_dim.h, last->output_dim.w, last->output_dim.h, 0, 0, tk::dnn::POOLING_AVERAGE); + new tk::dnn::Conv2d (&net, out_channel, 1, 1, 1, 1, 0, 0, ladder[li++], true); + + tk::dnn::Layer* act = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_SIGMOID); + new tk::dnn::Route(&net, &last, 1); + new tk::dnn::Shortcut(&net, act, true); + + //interpolate + new tk::dnn::Resize(&net, 1,2,2); + new tk::dnn::Shortcut(&net, down_out[1-i]); + + // //up-dense + new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true); + last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + up_out.push_back(last); + } + + + // for(int i=2;i>=0;--i){ + // new tk::dnn::Route(&net, &up_out[i], 1); + new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, conv_out[ci++], true); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); + new tk::dnn::Conv2d (&net, classes, 3, 3, 1, 1, 1, 1, conv_out[ci++], false); + /*up_out[i] =*/ new tk::dnn::Resize(&net, classes, net.input_dim.h, net.input_dim.w, true, tk::dnn::ResizeMode_t::LINEAR); + // } + + new tk::dnn::Softmax(&net); + + const char *output_bin = "shelfnet_mapillary/debug/softmax.bin"; + + // Load input + dnnType *data; + dnnType *input_h; + readBinaryFile(input_bin, dim.tot(), &input_h, &data); + std::cout<<"Input:"< 64) FatalError("Batch dim not supported"); - std::string net_name; - removePathAndExtension(net, net_name); - bool mapillary_15_colormap = false; - if( n_classes == 15 && - ( net_name == "shelfnet_mapillary_fp32" || - net_name == "shelfnet_mapillary_fp16" || - net_name == "shelfnet_mapillary_int8" ) ) - mapillary_15_colormap = true; - //net initialization tk::dnn::SegmentationNN segNN; segNN.init(net, n_classes, n_batch); @@ -130,7 +121,7 @@ int main(int argc, char *argv[]) { width = frame.cols; //inference - segNN.updateOriginal(frame, true, mapillary_15_colormap); + segNN.updateOriginal(frame, true); if(show) segNN.draw(); diff --git a/docs/README_seg.md b/docs/README_seg.md index 367a3a6..2479b55 100644 --- a/docs/README_seg.md +++ b/docs/README_seg.md @@ -39,14 +39,11 @@ where * `````` if set to 0 the demo will not show the visualization but save the video into result.mp4 (if n-batches ==1) * `````` if set to 0 (default) the demo will run, otherwise the evaluation of a dataset will run and the output of the segmentation will be saved. Attention: this is under development and paths are embedded, so change them in the code in advance. -N.b. By default it is used FP32 inference - -
+NB) By default it is used FP32 inference +NB) The batching is not used to work on more streams, rather to work on more tiles of the same image. Shelfnet never resized the input image, therefore for images greater than 1024x1024 tiles of 1024x1024 are given in input to the network in batch. ![gif](output.gif "Results on yolo_test.mp4") -
- For other demo videos refer to [this playlist](https://www.youtube.com/playlist?list=PLv0nEQYDD45y5EdSiywwCGPBmJVUzIWwe). @@ -60,6 +57,22 @@ For other demo videos refer to [this playlist](https://www.youtube.com/playlist? 1. Zhuang, Juntang, et al. "ShelfNet for fast semantic segmentation." Proceedings of the IEEE International Conference on Computer Vision Workshops. 2019. +## FPS Results + +Inference FPS of shelfnet with tkDNN, average of 1200 images on + * RTX 2080Ti (CUDA 10.2, TensorRT 7.0.0, Cudnn 7.6.5); + +| Platform | Test | Phase | FP32, ms | FP32, FPS | FP16, ms | FP16, FPS | INT8, ms | INT8, FPS | +| :------: | :-----: | :-----: | :-----: | :-----: | :-----: | :-----: | :-----: | :-----: | +| RTX 2080Ti | shelfnet 1024x1024 (B=1) | pre | 6.11863 | 163.435 | 5.81465 | 171.979 | 5.88699 | 169.866 | +| RTX 2080Ti | shelfnet 1024x1024 (B=1) | inf | 11.5464 | 86.6074 | 7.35396 | 135.981 | 6.37623 | 156.832 | +| RTX 2080Ti | shelfnet 1024x1024 (B=1) | post | 4.09058 | 244.464 | 3.91961 | 255.128 | 4.07343 | 245.493 | +| RTX 2080Ti | shelfnet 1024x1024 (B=1) | tot | 21.7556 | 45.9652 | 17.0882 | 58.5199 | 16.3366 | 61.2121 | +| RTX 2080Ti | shelfnet 2048x2048 (B=4) | pre | 25.435 | 39.3158 | 25.2953 | 39.5331 | 25.9303 | 38.565 | +| RTX 2080Ti | shelfnet 2048x2048 (B=4) | inf | 36.5015 | 27.3961 | 17.0534 | 58.6395 | 15.6061 | 64.0773 | +| RTX 2080Ti | shelfnet 2048x2048 (B=4) | post | 17.3917 | 57.4985 | 17.1649 | 58.2583 | 17.5539 | 56.9675 | +| RTX 2080Ti | shelfnet 2048x2048 (B=4) | tot | 79.3283 | 12.6058 | 59.5136 | 16.8029 | 59.0903 | 16.9233 | + ## Known issues When creating the rt file all the checks returns errors. It is due to a different resize function and handling of the original ShelfNet outputs. diff --git a/include/tkDNN/NetworkViz.h b/include/tkDNN/NetworkViz.h index 514fefc..2cf8009 100644 --- a/include/tkDNN/NetworkViz.h +++ b/include/tkDNN/NetworkViz.h @@ -5,8 +5,8 @@ namespace tk { namespace dnn { -cv::Mat vizFloat2colorMap(cv::Mat map, double min=0, double max=0, bool mapillary_15=false); -cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int img_h, int img_w, double min=0, double max=0, bool mapillary_15=false); +cv::Mat vizFloat2colorMap(cv::Mat map, double min=0, double max=0, int classes=19); +cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int img_h, int img_w, double min=0, double max=0, int classes=19); cv::Mat vizLayer2Mat(tk::dnn::Network *net, int layer, int imgdim = 1000); }} diff --git a/include/tkDNN/SegmentationNN.h b/include/tkDNN/SegmentationNN.h index d50bbc4..b691cbc 100644 --- a/include/tkDNN/SegmentationNN.h +++ b/include/tkDNN/SegmentationNN.h @@ -97,7 +97,7 @@ class SegmentationNN { * * @param bi batch index */ - void postprocess(const int bi=0, bool appy_colormap = true, bool mapillary_15=false) { + void postprocess(const int bi=0, bool appy_colormap = true) { dnnType *rt_out = (dnnType *)netRT->buffersRT[1]+ netRT->buffersDIM[1].tot()*bi; dataDim_t odim = netRT->output_dim; @@ -112,7 +112,7 @@ class SegmentationNN { cv::Mat colored; if(appy_colormap) - colored = vizData2Mat(tmpOutData_h, vdim, netRT->input_dim.h, netRT->input_dim.w, 0, classes, mapillary_15); + colored = vizData2Mat(tmpOutData_h, vdim, netRT->input_dim.h, netRT->input_dim.w, 0, classes, classes); else{ cv::Mat colored_fp32 (cv::Size(odim.w, odim.h),CV_32FC1, tmpOutData_h); colored_fp32.convertTo(colored, CV_8UC1); @@ -234,7 +234,7 @@ class SegmentationNN { } } - void updateOriginal(cv::Mat frame, bool apply_colormap=true, bool mapillary_15=false){ + void updateOriginal(cv::Mat frame, bool apply_colormap=true){ std::vector splitted_frames; int H, W, net_H, net_W; @@ -347,7 +347,7 @@ class SegmentationNN { cv::Mat colored; if(apply_colormap) - colored = vizData2Mat(tmpOutData_h, vdim, netRT->input_dim.h, netRT->input_dim.w, 0, classes, mapillary_15); + colored = vizData2Mat(tmpOutData_h, vdim, netRT->input_dim.h, netRT->input_dim.w, 0, classes, classes); else{ cv::Mat colored_fp32 (cv::Size(odim.w, odim.h),CV_32FC1, tmpOutData_h); colored_fp32.convertTo(colored, CV_8UC1); diff --git a/src/NetworkViz.cpp b/src/NetworkViz.cpp index aaf5d55..b6c95d9 100644 --- a/src/NetworkViz.cpp +++ b/src/NetworkViz.cpp @@ -12,62 +12,77 @@ cv::Mat mapillary_15_map(cv::Mat adjMap){ // cv::waitKey(0); cv::Mat M1(1, 256, CV_8UC1), M2(1, 256, CV_8UC1), M3(1, 256, CV_8UC1); + //animal M3.at(0)=165; M2.at(0)=42; M1.at(0)=45; + //curb M3.at(1)=196; M2.at(1)=196; M1.at(1)=196; + //barrier M3.at(2)=90; M2.at(2)=120; M1.at(2)=150; + //road M3.at(3)=128; M2.at(3)=64; M1.at(3)=128; + //building M3.at(4)=70; M2.at(4)=70; M1.at(4)=70; + //person M3.at(5)=220; M2.at(5)=20; M1.at(5)=60; + //roadmark M3.at(6)=255; M2.at(6)=255; M1.at(6)=255; + //nature M3.at(7)=107; M2.at(7)=142; M1.at(7)=35; + //sky M3.at(8)=70; M2.at(8)=130; M1.at(8)=180; + //billboard M3.at(9)=220; M2.at(9)=220; M1.at(9)=220; + //pole M3.at(10)=153; M2.at(10)=153; M1.at(10)=153; + //traffic sign M3.at(11)=128; M2.at(11)=128; M1.at(11)=128; + //bike M3.at(12)=119; M2.at(12)=11; M1.at(12)=32; + //vehicle M3.at(13)=0; M2.at(13)=0; M1.at(13)=142; + //void for(int i=14;i<256;i++) { M1.at(i)=0; @@ -93,7 +108,256 @@ cv::Mat mapillary_15_map(cv::Mat adjMap){ } -cv::Mat vizFloat2colorMap(cv::Mat map,double min, double max, bool mapillary_15) { +cv::Mat berkeley_20_map(cv::Mat adjMap){ + + cv::Mat M1(1, 256, CV_8UC1), M2(1, 256, CV_8UC1), M3(1, 256, CV_8UC1); + + //road + M3.at(0)=128; + M2.at(0)=64; + M1.at(0)=128; + + //sidewalk + M3.at(1)=244; + M2.at(1)=35; + M1.at(1)=232; + + //building + M3.at(2)=70; + M2.at(2)=70; + M1.at(2)=70; + + //wall + M3.at(3)=102; + M2.at(3)=102; + M1.at(3)=156; + + //fence + M3.at(4)=90; + M2.at(4)=120; + M1.at(4)=150; + + //pole + M3.at(5)=153; + M2.at(5)=153; + M1.at(5)=153; + + //traffic light + M3.at(6)=250; + M2.at(6)=170; + M1.at(6)=30; + + //traffic sign + M3.at(7)=128; + M2.at(7)=128; + M1.at(7)=128; + + //nature + M3.at(8)=107; + M2.at(8)=142; + M1.at(8)=35; + + //ground + M3.at(9)=0; + M2.at(9)=192; + M1.at(9)=0; + + //sky + M3.at(10)=70; + M2.at(10)=130; + M1.at(10)=180; + + //person + M3.at(11)=220; + M2.at(11)=20; + M1.at(11)=60; + + //rider + M3.at(12)=255; + M2.at(12)=0; + M1.at(12)=100; + + //car + M3.at(13)=0; + M2.at(13)=0; + M1.at(13)=142; + + //truck + M3.at(14)=0; + M2.at(14)=0; + M1.at(14)=70; + + //bus + M3.at(15)=0; + M2.at(15)=60; + M1.at(15)=100; + + //train + M3.at(16)=0; + M2.at(16)=0; + M1.at(16)=192; + + //motorbike + M3.at(17)=0; + M2.at(17)=0; + M1.at(17)=230; + + //bike + M3.at(18)=119; + M2.at(18)=11; + M1.at(18)=32; + + //void + for(int i=19;i<256;i++) + { + M1.at(i)=0; + M2.at(i)=0; + M3.at(i)=0; + } + + cv::Mat r1,r2,r3; + + cv::LUT(adjMap,M1,r1); + cv::LUT(adjMap,M2,r2); + cv::LUT(adjMap,M3,r3); + + std::vector planes; + planes.push_back(r1); + planes.push_back(r2); + planes.push_back(r3); + + cv::Mat dst; + cv::merge(planes,dst); + return dst; + +} + +cv::Mat cityscapes_19_map(cv::Mat adjMap){ + + cv::Mat M1(1, 256, CV_8UC1), M2(1, 256, CV_8UC1), M3(1, 256, CV_8UC1); + + //road + M3.at(0)=128; + M2.at(0)=64; + M1.at(0)=128; + + //sidewalk + M3.at(1)=244; + M2.at(1)=35; + M1.at(1)=232; + + //building + M3.at(2)=70; + M2.at(2)=70; + M1.at(2)=70; + + //wall + M3.at(3)=102; + M2.at(3)=102; + M1.at(3)=156; + + //fence + M3.at(4)=190; + M2.at(4)=153; + M1.at(4)=153; + + //pole + M3.at(5)=153; + M2.at(5)=153; + M1.at(5)=153; + + //traffic light + M3.at(6)=250; + M2.at(6)=170; + M1.at(6)=30; + + //traffic sign + M3.at(7)=220; + M2.at(7)=220; + M1.at(7)=0; + + //vegetation + M3.at(8)=107; + M2.at(8)=142; + M1.at(8)=35; + + //terrain + M3.at(9)=152; + M2.at(9)=251; + M1.at(9)=152; + + //sky + M3.at(10)=70; + M2.at(10)=130; + M1.at(10)=180; + + //person + M3.at(11)=220; + M2.at(11)=20; + M1.at(11)=60; + + //rider + M3.at(12)=255; + M2.at(12)=0; + M1.at(12)=0; + + //car + M3.at(13)=0; + M2.at(13)=0; + M1.at(13)=142; + + //truck + M3.at(14)=0; + M2.at(14)=0; + M1.at(14)=70; + + //bus + M3.at(15)=0; + M2.at(15)=60; + M1.at(15)=100; + + //train + M3.at(16)=0; + M2.at(16)=80; + M1.at(16)=100; + + //motorcycle + M3.at(17)=0; + M2.at(17)=0; + M1.at(17)=230; + + //bicycle + M3.at(18)=119; + M2.at(18)=11; + M1.at(18)=32; + + //void + for(int i=19;i<256;i++) + { + M1.at(i)=0; + M2.at(i)=0; + M3.at(i)=0; + } + + cv::Mat r1,r2,r3; + + cv::LUT(adjMap,M1,r1); + cv::LUT(adjMap,M2,r2); + cv::LUT(adjMap,M3,r3); + + std::vector planes; + planes.push_back(r1); + planes.push_back(r2); + planes.push_back(r3); + + cv::Mat dst; + cv::merge(planes,dst); + return dst; + +} + + +cv::Mat vizFloat2colorMap(cv::Mat map,double min, double max, int classes) { if(min == 0 && max == 0) cv::minMaxIdx(map, &min, &max); @@ -101,11 +365,22 @@ cv::Mat vizFloat2colorMap(cv::Mat map,double min, double max, bool mapillary_15) cv::Mat adjMap; cv::Mat falseColorsMap; - if(mapillary_15){ + switch (classes) + { + case 15: map.convertTo(adjMap,CV_8UC1); falseColorsMap = mapillary_15_map(adjMap); - } - else{ + break; + case 20: + map.convertTo(adjMap,CV_8UC1); + falseColorsMap = berkeley_20_map(adjMap); + break; + case 19: + map.convertTo(adjMap,CV_8UC1); + falseColorsMap = cityscapes_19_map(adjMap); + break; + + default: // expand your range to 0..255. Similar to histEq(); map.convertTo(adjMap,CV_8UC1, 255 / (max-min), -min); applyColorMap(adjMap, falseColorsMap, cv::COLORMAP_JET); @@ -113,7 +388,7 @@ cv::Mat vizFloat2colorMap(cv::Mat map,double min, double max, bool mapillary_15) return falseColorsMap; } -cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int img_h, int img_w, double min, double max, bool mapillary_15) { +cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int img_h, int img_w, double min, double max, int classes) { dnnType *data = nullptr; // copy to CPU @@ -129,7 +404,7 @@ cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int img_h, int i cv::Mat grid = cv::Mat(gridSize, CV_8UC3, cv::Scalar(0)); for(int i=0; i Date: Tue, 20 Jul 2021 17:30:32 +0200 Subject: [PATCH 17/17] Update READMEs Signed-off-by: Micaela Verucchi --- README.md | 25 +++++++++++++++---------- docs/README_seg.md | 12 +++++++++++- 2 files changed, 26 insertions(+), 11 deletions(-) diff --git a/README.md b/README.md index 4258456..7dcd080 100644 --- a/README.md +++ b/README.md @@ -18,7 +18,7 @@ If you use tkDNN in your research, please cite the [following paper](https://iee ``` ### What's new (20 July 2021) -- [x] Support to sematic segmentation [REAME](docs/README_seg.md) +- [x] Support to sematic segmentation [README](docs/README_seg.md) - [ ] Support to TensorRT8 (WIP) ## FPS Results @@ -93,16 +93,21 @@ Results for COCO val 2017 (5k images), on RTX 2080Ti, with conf threshold=0.001 - [Known issues with tkDNN on Windows](#known-issues-with-tkdnn-on-windows) - - - ## Dependencies -This branch works on every NVIDIA GPU that supports the dependencies: -* CUDA 10.0 -* CUDNN 7.603 -* TENSORRT 6.01 -* OPENCV 3.4 -* yaml-cpp 0.5.2 (sudo apt install libyaml-cpp-dev) +This branch works on every NVIDIA GPU that supports the following (latest tested) dependencies: +* CUDA 11.0 (or >= 10) +* cuDNN 8.0.4 (or >= 7.3) +* TensorRT 7.2.0 (or >=5) +* OpenCV 4.5.2 (or >=4) +* cmake 3.21 (or >= 3.15) +* yaml-cpp 0.5.2 +* eigen3 3.3.4 +* curl 7.58 + +``` +sudo apt install libyaml-cpp-dev curl libeigen3-dev + +``` ## About OpenCV To compile and install OpenCV4 with contrib us the script ```install_OpenCV4.sh```. It will download and compile OpenCV in Download folder. diff --git a/docs/README_seg.md b/docs/README_seg.md index 2479b55..3bd1060 100644 --- a/docs/README_seg.md +++ b/docs/README_seg.md @@ -59,8 +59,9 @@ For other demo videos refer to [this playlist](https://www.youtube.com/playlist? ## FPS Results -Inference FPS of shelfnet with tkDNN, average of 1200 images on +Inference FPS of shelfnet with tkDNN, average of 1200 images on: * RTX 2080Ti (CUDA 10.2, TensorRT 7.0.0, Cudnn 7.6.5); + * Xavier AGX, Jetpack 4.3 (CUDA 10.0, CUDNN 7.6.3, tensorrt 6.0.1 ); | Platform | Test | Phase | FP32, ms | FP32, FPS | FP16, ms | FP16, FPS | INT8, ms | INT8, FPS | | :------: | :-----: | :-----: | :-----: | :-----: | :-----: | :-----: | :-----: | :-----: | @@ -72,6 +73,15 @@ Inference FPS of shelfnet with tkDNN, average of 1200 images on | RTX 2080Ti | shelfnet 2048x2048 (B=4) | inf | 36.5015 | 27.3961 | 17.0534 | 58.6395 | 15.6061 | 64.0773 | | RTX 2080Ti | shelfnet 2048x2048 (B=4) | post | 17.3917 | 57.4985 | 17.1649 | 58.2583 | 17.5539 | 56.9675 | | RTX 2080Ti | shelfnet 2048x2048 (B=4) | tot | 79.3283 | 12.6058 | 59.5136 | 16.8029 | 59.0903 | 16.9233 | +| AGX Xavier | shelfnet 1024x1024 (B=1) | pre | 8.0174 | 124.729 | 7.5117 | 133.126 | 7.47333 | 133.809 | +| AGX Xavier | shelfnet 1024x1024 (B=1) | inf | 72.4173 | 13.8089 | 37.505 | 26.6631 | 31.3286 | 31.9197 | +| AGX Xavier | shelfnet 1024x1024 (B=1) | post | 8.89958 | 112.365 | 8.83576 | 113.176 | 9.42655 | 106.083 | +| AGX Xavier | shelfnet 1024x1024 (B=1) | tot | 89.3342 | 11.1939 | 53.8525 | 18.5692 | 48.2285 | 20.7346 | +| AGX Xavier | shelfnet 2048x2048 (B=4) | pre | 47.1454 | 21.211 | 21.6475 | 46.1947 | 21.4201 | 46.6851 | +| AGX Xavier | shelfnet 2048x2048 (B=4) | inf | 266.537 | 3.75183 | 128.321 | 7.79293 | 107.621 | 9.29185 | +| AGX Xavier | shelfnet 2048x2048 (B=4) | post | 44.0711 | 22.6906 | 40.1732 | 24.8922 | 39.873 | 25.0796 | +| AGX Xavier | shelfnet 2048x2048 (B=4) | tot | 357.753 | 2.79522 | 190.142 | 5.25922 | 168.914 | 5.92016 | + ## Known issues