From 5595b8037b2ee3b747e456244b0c63943a38938e Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Tue, 22 Aug 2017 01:32:59 -0700 Subject: [PATCH] interpret --- CMakeLists.txt | 7 ++- include/Layer.h | 30 +++++++++--- src/NetworkRT.cpp | 10 ++-- src/Region.cpp | 92 ++++++++++++++++++++++------------- src/pluginsRT/RegionRT.cpp | 5 +- src/utils.cpp | 3 +- tests/yolo-tiny/yolo-tiny.cpp | 6 +-- tests/yolo/yolo.cpp | 13 +++-- 8 files changed, 106 insertions(+), 60 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index d206dca..2c4497c 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -61,9 +61,14 @@ add_executable(test_yolo_tiny tests/yolo-tiny/yolo-tiny.cpp) target_link_libraries(test_yolo_tiny tkDNN) add_executable(test_rtinference tests/test_rtinference/rtinference.cpp) -target_link_libraries(test_rtinference tkDNN_static) +target_link_libraries(test_rtinference tkDNN) #install +if (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT) + set (CMAKE_INSTALL_PREFIX "${CMAKE_BINARY_DIR}/install" + CACHE PATH "default install path" FORCE) +endif() +message("install dir:" ${CMAKE_INSTALL_PREFIX}) install(DIRECTORY include/ DESTINATION include/${CMAKE_PROJECT_NAME} FILES_MATCHING PATTERN "*.h") install(TARGETS tkDNN DESTINATION lib) diff --git a/include/Layer.h b/include/Layer.h index aa7d7ca..3cc2756 100644 --- a/include/Layer.h +++ b/include/Layer.h @@ -276,6 +276,11 @@ public: struct box { float x, y, w, h; }; +struct sortable_bbox { + int index; + int cl; + float **probs; +}; /** Region layer @@ -284,29 +289,42 @@ struct box { class Region : public Layer { public: - Region(Network *net, int classes, int coords, int num, float thresh, const char* fname_weights); + Region(Network *net, int classes, int coords, int num); virtual ~Region(); virtual layerType_t getLayerType() { return LAYER_REGION; }; + int classes, coords, num; + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); +}; - dnnType *bias_h, *bias_d; +class RegionInterpret { + +public: + RegionInterpret(dataDim_t input_dim, dataDim_t output_dim, + int classes, int coords, int num, float thresh, const char* fname_weights); + ~RegionInterpret(); + + dataDim_t input_dim, output_dim; + dnnType *bias_h, *bias_d; //anchors int classes, coords, num; float thresh; + + + box *boxes; + float **probs; + sortable_bbox *s; box res_boxes[256]; int res_boxes_n; - int entry_index(int batch, int location, int entry); box get_region_box(float *x, float *biases, int n, int index, int i, int j, int w, int h, int stride); void get_region_boxes( float *input, int w, int h, int netw, int neth, float thresh, float **probs, box *boxes, int only_objectness, int *map, float tree_thresh, int relative); void correct_region_boxes(box *boxes, int n, int w, int h, int netw, int neth, int relative); - void interpretData(); + void interpretData(dnnType *data_h); void showImageResult(dnnType *input_h); }; - - } #endif //LAYER_H diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 3eb486a..8122810 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -45,7 +45,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) { builderRT->setMaxBatchSize(1); builderRT->setMaxWorkspaceSize(1 << 30); - +/* //change datatype based on system specs if(builderRT->platformHasFastInt8()) { BatchStream bstream({32,dim.c, dim.h, dim.w}, 32, 1); @@ -57,7 +57,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) { dtRT = DataType::kHALF; builderRT->setHalf2Mode(true); } - +*/ //add input layer ITensor *input = networkRT->addInput("data", dtRT, DimsCHW{ dim.c, dim.h, dim.w}); @@ -298,7 +298,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Region *l) { //std::cout<<"convert Region\n"; //std::cout<<"New plugin REGION\n"; - IPlugin *plugin = new RegionRT(l->classes, l->coords, l->num, l->thresh); + IPlugin *plugin = new RegionRT(l->classes, l->coords, l->num); IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); checkNULL(lRT); return lRT; @@ -338,8 +338,8 @@ public: if(name.find("Region") == 0) { RegionRT *r = new RegionRT(readBUF(buf), //classes readBUF(buf), //coords - readBUF(buf), //num - readBUF(buf)); //thesh + readBUF(buf)); //num + r->c = readBUF(buf); r->h = readBUF(buf); r->w = readBUF(buf); diff --git a/src/Region.cpp b/src/Region.cpp index b28ee47..73c2778 100644 --- a/src/Region.cpp +++ b/src/Region.cpp @@ -10,15 +10,13 @@ namespace tkDNN { -Region::Region(Network *net, int classes, int coords, int num, float thresh, const char* fname_weights) : +Region::Region(Network *net, int classes, int coords, int num) : Layer(net) { this->classes = classes; this->coords = coords; this->num = num; - this->thresh = thresh; - this->res_boxes_n = 0; - + // same output_dim.n = input_dim.n; output_dim.c = input_dim.c; @@ -26,9 +24,6 @@ Region::Region(Network *net, int classes, int coords, int num, float thresh, con output_dim.w = input_dim.w; output_dim.l = input_dim.l; - //load anchors - readBinaryFile(fname_weights, 2*num, &bias_h, &bias_d); - checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) ); } @@ -36,10 +31,12 @@ Region::~Region() { checkCuda( cudaFree(dstData) ); } -int Region::entry_index(int batch, int location, int entry) { +int entry_index(int batch, int location, int entry, + int coords, int classes, dataDim_t &input_dim, dataDim_t &output_dim) { int n = location / (input_dim.w*input_dim.h); int loc = location % (input_dim.w*input_dim.h); - return batch*output_dim.tot() + n*input_dim.w*input_dim.h*(coords+classes+1) + entry*input_dim.w*input_dim.h + loc; + return batch*output_dim.tot() + n*input_dim.w*input_dim.h*(coords+classes+1) + + entry*input_dim.w*input_dim.h + loc; } dnnType* Region::infer(dataDim_t &dim, dnnType* srcData) { @@ -48,16 +45,16 @@ dnnType* Region::infer(dataDim_t &dim, dnnType* srcData) { for (int b = 0; b < dim.n; ++b){ for(int n = 0; n < num; ++n){ - int index = entry_index(b, n*dim.w*dim.h, 0); + int index = entry_index(b, n*dim.w*dim.h, 0, coords, classes, input_dim, output_dim); activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h); - index = entry_index(b, n*dim.w*dim.h, coords); + index = entry_index(b, n*dim.w*dim.h, coords, coords, classes, input_dim, output_dim); activationLOGISTICForward(srcData + index, dstData + index, dim.w*dim.h); } } //softmax start - int index = entry_index(0, 0, coords + 1); + int index = entry_index(0, 0, coords + 1, coords, classes, input_dim, output_dim); softmaxForward(srcData + index, classes, output_dim.n*num, output_dim.tot()/num, output_dim.w*output_dim.h, 1, output_dim.w*output_dim.h, 1, dstData + index); @@ -66,7 +63,42 @@ dnnType* Region::infer(dataDim_t &dim, dnnType* srcData) { } -box Region::get_region_box(float *x, float *biases, int n, int index, int i, int j, int w, int h, int stride) +/* Intepret class */ +RegionInterpret::RegionInterpret(dataDim_t input_dim, dataDim_t output_dim, + int classes, int coords, int num, float thresh, const char* fname_weights) { + + this->input_dim = input_dim; + this->output_dim = output_dim; + + this->classes = classes; + this->coords = coords; + this->num = num; + this->thresh = thresh; + this->res_boxes_n = 0; + + int tot = output_dim.w*output_dim.h*num; + boxes = (box*) malloc(tot*sizeof(box)); + probs = (float**) malloc(tot*sizeof(float *)); + for(int j = 0; j < tot; ++j) probs[j] = (float*) malloc((classes + 1)*sizeof(float *)); + s = (sortable_bbox*) malloc(tot*sizeof(sortable_bbox)); + + //load anchors + readBinaryFile(fname_weights, 2*num, &bias_h, &bias_d); +} + +RegionInterpret::~RegionInterpret() { + + delete [] boxes; + for(int j = 0; j < output_dim.w*output_dim.h*num; ++j) + delete [] probs[j]; + delete [] probs; + delete [] s; + + delete [] bias_h; + checkCuda( cudaFree(bias_d) ); +} + +box RegionInterpret::get_region_box(float *x, float *biases, int n, int index, int i, int j, int w, int h, int stride) { box b; b.x = (i + x[index + 0*stride]) / w; @@ -76,7 +108,7 @@ box Region::get_region_box(float *x, float *biases, int n, int index, int i, int return b; } -void Region::get_region_boxes( float *input, int w, int h, int netw, int neth, float thresh, +void RegionInterpret::get_region_boxes( float *input, int w, int h, int netw, int neth, float thresh, float **probs, box *boxes, int only_objectness, int *map, float tree_thresh, int relative) { int lh = output_dim.h; @@ -92,14 +124,17 @@ void Region::get_region_boxes( float *input, int w, int h, int netw, int neth, for(int j = 0; j < classes; ++j){ probs[index][j] = 0; } - int obj_index = entry_index(0, n*lw*lh + i, coords); - int box_index = entry_index(0, n*lw*lh + i, 0); + int obj_index = entry_index(0, n*lw*lh + i, + coords, coords, classes, output_dim, output_dim); + int box_index = entry_index(0, n*lw*lh + i, 0, + coords, classes, output_dim, output_dim); float scale = predictions[obj_index]; boxes[index] = get_region_box(predictions, bias_h, n, box_index, col, row, lw, lh, lw*lh); float max = 0; for(int j = 0; j < classes; ++j){ - int class_index = entry_index(0, n*lw*lh + i, coords + 1 + j); + int class_index = entry_index(0, n*lw*lh + i, coords + 1 + j, + coords, classes, output_dim, output_dim); float prob = scale*predictions[class_index]; probs[index][j] = (prob > thresh) ? prob : 0; if(prob > max) max = prob; @@ -111,7 +146,7 @@ void Region::get_region_boxes( float *input, int w, int h, int netw, int neth, } -void Region::correct_region_boxes(box *boxes, int n, int w, int h, int netw, int neth, int relative) { +void RegionInterpret::correct_region_boxes(box *boxes, int n, int w, int h, int netw, int neth, int relative) { int i; int new_w=0; int new_h=0; @@ -142,11 +177,6 @@ void Region::correct_region_boxes(box *boxes, int n, int w, int h, int netw, int //############################ BOX PROBABILITY UTILS ############################ -struct sortable_bbox { - int index; - int cl; - float **probs; -}; int nms_comparator(const void *pa, const void *pb) { sortable_bbox a = *(sortable_bbox *)pa; sortable_bbox b = *(sortable_bbox *)pb; @@ -196,21 +226,15 @@ int max_index(float *a, int n) { -void Region::interpretData() { +void RegionInterpret::interpretData(dnnType *data_h) { - int imW = net->input_dim.w, imH = net->input_dim.h; + int imW = input_dim.w, imH = input_dim.h; int tot = output_dim.w*output_dim.h*num; - float *lel = new dnnType[output_dim.tot()]; - cudaMemcpy(lel, dstData, output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost); - box *boxes = (box*) calloc(tot, sizeof(box)); - float **probs = (float**) calloc(tot, sizeof(float *)); - for(int j = 0; j < tot; ++j) probs[j] = (float*)calloc(classes + 1, sizeof(float *)); - get_region_boxes(lel, imW, imH, output_dim.w, output_dim.h, thresh, probs, boxes, 0, 0, 0.5, 1); + get_region_boxes(data_h, imW, imH, output_dim.w, output_dim.h, thresh, probs, boxes, 0, 0, 0.5, 1); //delete repeats - sortable_bbox *s = (sortable_bbox*)calloc(tot, sizeof(sortable_bbox)); for(int i = 0; i < tot; ++i){ s[i].index = i; s[i].cl = classes; @@ -230,13 +254,13 @@ void Region::interpretData() { } } } - free(s); // //print results for(int i = 0; i < tot; ++i){ int cl = max_index(probs[i], classes); float prob = probs[i][cl]; + if(prob > thresh) { box b = boxes[i]; int x = (b.x-b.w/2.)*imW; @@ -255,7 +279,7 @@ void Region::interpretData() { } } -void Region::showImageResult(dnnType *input_h) { +void RegionInterpret::showImageResult(dnnType *input_h) { #ifdef OPENCV dataDim_t dim = net->input_dim; diff --git a/src/pluginsRT/RegionRT.cpp b/src/pluginsRT/RegionRT.cpp index 880701b..f1e7e9e 100644 --- a/src/pluginsRT/RegionRT.cpp +++ b/src/pluginsRT/RegionRT.cpp @@ -4,12 +4,11 @@ class RegionRT : public IPlugin { public: - RegionRT(int classes, int coords, int num, float thresh) { + RegionRT(int classes, int coords, int num) { this->classes = classes; this->coords = coords; this->num = num; - this->thresh = thresh; } ~RegionRT(){ @@ -78,7 +77,6 @@ public: tkDNN::writeBUF(buf, classes); tkDNN::writeBUF(buf, coords); tkDNN::writeBUF(buf, num); - tkDNN::writeBUF(buf, thresh); tkDNN::writeBUF(buf, c); tkDNN::writeBUF(buf, h); tkDNN::writeBUF(buf, w); @@ -86,7 +84,6 @@ public: int c, h, w; int classes, coords, num; - float thresh; int entry_index(int batch, int location, int entry, int batchSize) { int n = location / (w*h); diff --git a/src/utils.cpp b/src/utils.cpp index 0240874..253833f 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -86,7 +86,8 @@ int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) { int diffs = 0; for(int i=0; i eps) { + if(data_h[i] != data_h[i] && correct_h[i] != correct_h[i] && //nan control + fabs(data_h[i] - correct_h[i]) > eps) { diffs += 1; if(diffs == 1) std::cout<<"\n"; diff --git a/tests/yolo-tiny/yolo-tiny.cpp b/tests/yolo-tiny/yolo-tiny.cpp index 954fe5f..465c601 100644 --- a/tests/yolo-tiny/yolo-tiny.cpp +++ b/tests/yolo-tiny/yolo-tiny.cpp @@ -49,7 +49,7 @@ int main() { tkDNN::Conv2d c12(&net, 512, 3, 3, 1, 1, 1, 1, c12_bin, true); tkDNN::Activation a12(&net, tkDNN::ACTIVATION_LEAKY); tkDNN::Conv2d c13(&net, 425, 1, 1, 1, 1, 0, 0, c13_bin, false); - tkDNN::Region g14(&net, 80, 4, 5, 0.6f, g14_bin); + tkDNN::Region g14(&net, 80, 4, 5); // Load input dnnType *data; @@ -89,9 +89,5 @@ int main() { std::cout<<"CUDNN vs correct"; checkResult(out_dim, out_data, out); std::cout<<"TRT vs correct"; checkResult(out_dim, out_data2, out); std::cout<<"CUDNN vs TRT "; checkResult(out_dim, out_data, out_data2); - - std::cout<<"\n\nDetected objects: \n"; - g14.interpretData(); - g14.showImageResult(input_h); return 0; } diff --git a/tests/yolo/yolo.cpp b/tests/yolo/yolo.cpp index 50aec86..a485542 100644 --- a/tests/yolo/yolo.cpp +++ b/tests/yolo/yolo.cpp @@ -97,7 +97,9 @@ int main() { tkDNN::Conv2d c29(&net, 1024, 3, 3, 1, 1, 1, 1, c29_bin, true); tkDNN::Activation a29(&net, tkDNN::ACTIVATION_LEAKY); tkDNN::Conv2d c30(&net, 425, 1, 1, 1, 1, 0, 0, c30_bin, false); - tkDNN::Region g31(&net, 80, 4, 5, 0.6f, g31_bin); + tkDNN::Region g31(&net, 80, 4, 5); + + tkDNN::RegionInterpret rI(dim, g31.output_dim, 80, 4, 5, 0.6f, g31_bin); // Load input dnnType *data; @@ -120,7 +122,7 @@ int main() { TIMER_STOP dim1.print(); } - + tkDNN::dataDim_t dim2 = dim; printCenteredTitle(" TENSORRT inference ", '=', 30); { dim2.print(); @@ -139,7 +141,10 @@ int main() { std::cout<<"CUDNN vs TRT "; checkResult(out_dim, out_data, out_data2); std::cout<<"\n\nDetected objects: \n"; - g31.interpretData(); - g31.showImageResult(input_h); + dnnType *output_h = new dnnType[rI.output_dim.tot()]; + checkCuda(cudaMemcpy(output_h, out_data2, + rI.output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost)); + rI.interpretData(output_h); + rI.showImageResult(input_h); return 0; }