diff --git a/CMakeLists.txt b/CMakeLists.txt index 8043e4e..021189b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -27,7 +27,7 @@ if(NVINFER STREQUAL "NVINFER-NOTFOUND") endif() # compile -file(GLOB tkdnn_CUSRC "src/kernels/*.cu") +file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/*.cu") cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${NVINFER_INCLUDES}) cuda_add_library(kernels SHARED ${tkdnn_CUSRC}) diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index 295b2af..81ca865 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -9,7 +9,8 @@ #include #include -#include "Yolo3Detection.h" +// #include "Yolo3Detection.h" +#include "CenternetDetection.h" bool gRun; bool SAVE_RESULT = false; @@ -25,14 +26,15 @@ int main(int argc, char *argv[]) { signal(SIGINT, sig_handler); - char *net = "yolo3.rt"; + char *net = "resnet101_cnet.rt"; if(argc > 1) net = argv[1]; char *input = "../demo/yolo_test.mp4"; if(argc > 2) input = argv[2]; - tk::dnn::Yolo3Detection yolo; + // tk::dnn::Yolo3Detection yolo; + tk::dnn::CenternetDetection yolo; yolo.init(net); gRun = true; @@ -66,29 +68,31 @@ int main(int argc, char *argv[]) { // TODO: async infer yolo.update(dnn_input); - // draw dets - for(int i=0; iclassesNames[b.cl]; - float prob = b.prob; + frame = yolo.draw(dnn_input); + // // draw dets + // for(int i=0; iclassesNames[b.cl]; + // float prob = b.prob; - std::cout< +#include +#include +#include +#include +#include +#include + + +#include "tkdnn.h" + +void sort(dnnType *src_begin, dnnType *src_end, int *idsrc); +void topk(dnnType *src_begin, int *idsrc, int K, float *topk_scores, + int *topk_inds, float *topk_ys, float *topk_xs); +void sortAndTopKonDevice(dnnType *src_begin, int *idsrc, float *topk_scores, int *topk_inds, float *topk_ys, float *topk_xs, const int size, const int K, const int n_classes); +void subtractWithThreshold(dnnType *src_begin, dnnType *src_end, dnnType *src2_begin, dnnType *src_out); +void topKxyclasses(int *ids_begin, int *ids_end, const int K, const int size, const int wh, int *clses, int *xs, int *ys); +void topKxyAddOffset(int * ids_begin, const int K, const int size, int *intxs_begin, int *intys_begin, float *xs_begin, float *ys_begin, dnnType *src_begin); +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); diff --git a/include/tkDNN/CenternetDetection.h b/include/tkDNN/CenternetDetection.h new file mode 100644 index 0000000..8071112 --- /dev/null +++ b/include/tkDNN/CenternetDetection.h @@ -0,0 +1,112 @@ +#include +#include +#include +#include /* srand, rand */ +#include +#include +#include "utils.h" +#include +#include "kernels.h" +#include +#include // std::iota +#include // std::sort + + +#include +#include +#include + +#include "tkdnn.h" +#include "sorting.h" + +namespace tk { namespace dnn { + +/** + * + * @author Francesco Gatti + */ +class CenternetDetection { + + private: + tk::dnn::NetworkRT *netRT = nullptr; + dnnType *input_h, *input, *input_d; + + int ndets = 0; + // tk::dnn::Yolo::detection *dets = nullptr; + + cv::Mat imageF; + cv::Mat bgr[3]; + + // variable to test cnet on dog pictures + tk::dnn::dataDim_t dim; + tk::dnn::dataDim_t dim2; + cv::Size sz; + const char *input_bin = "../tests/resnet101_cnet/debug/input.bin"; + + // pre-process + tk::dnn::dataDim_t dim_hm; + tk::dnn::dataDim_t dim_wh; + tk::dnn::dataDim_t dim_reg; + float *topk_scores; + int *topk_inds_; + float *topk_ys_; + float *topk_xs_; + int *ids_d, *ids_, *ids_2, *ids_2d; + + float *scores, *scores_d; + int *clses, *clses_d; + int *topk_inds_d; + float *topk_ys_d; + float *topk_xs_d; + int *inttopk_xs_d, *inttopk_ys_d; + + + float *bbx0, *bby0, *bbx1, *bby1; + float *bbx0_d, *bby0_d, *bbx1_d, *bby1_d; + + float *target_coords; + + cv::Vec mean; + cv::Vec stddev; + cv::Mat src; + cv::Mat dst; + //processing + float toll = 0.000001; + int K = 100; + int width = 56; // TODO + + + public: + dnnType *rt_out[4]; + + float inp_height = 224;//512; + float inp_width = 224;//512; + + int classes = 80; + int num = 0; + int n_masks = 0; + float thresh = 0.3; + cv::Scalar colors[256]; + + // this is filled with results + std::vector detected; + // draw + std::vector coco_class_name; + + CenternetDetection() {} + + virtual ~CenternetDetection() {} + + /** + * Method used for inizialize the class + * + * @return Success of the initialization + */ + bool init(std::string tensor_path); + void testdog(); + cv::Mat draw(cv::Mat &frame); + void update(cv::Mat &frame); + +}; + +}} diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index f5ac6a7..470168a 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -292,7 +292,7 @@ public: Pooling(Network *net, int winH, int winW, int strideH, int strideW, int paddingH = 0, int paddingW = 0, - tkdnnPoolingMode_t pool_mode = POOLING_MAX); + tkdnnPoolingMode_t pool_mode = POOLING_MAX, bool final = false); virtual ~Pooling(); virtual layerType_t getLayerType() { return LAYER_POOLING; }; diff --git a/src/CenternetDetection.cpp b/src/CenternetDetection.cpp new file mode 100644 index 0000000..59bf827 --- /dev/null +++ b/src/CenternetDetection.cpp @@ -0,0 +1,617 @@ +#include "CenternetDetection.h" + +namespace tk { namespace dnn { + +float __colors[6][3] = { {1,0,1}, {0,0,1},{0,1,1},{0,1,0},{1,1,0},{1,0,0} }; +float get_color2(int c, int x, int max) +{ + float ratio = ((float)x/max)*5; + int i = floor(ratio); + int j = ceil(ratio); + ratio -= i; + float r = (1-ratio) * __colors[i % 6][c % 3] + ratio*__colors[j % 6][c % 3]; + //printf("%f\n", r); + return r; +} + +bool CenternetDetection::init(std::string tensor_path) { + std::cout<<(tensor_path).c_str()<<"\n"; + netRT = new tk::dnn::NetworkRT(NULL, (tensor_path).c_str() ); + + dim = tk::dnn::dataDim_t(1, 3, 224, 224, 1); + const char *coco_class_name_[] = { + "person", "bicycle", "car", "motorcycle", "airplane", + "bus", "train", "truck", "boat", "traffic light", "fire hydrant", + "stop sign", "parking meter", "bench", "bird", "cat", "dog", "horse", + "sheep", "cow", "elephant", "bear", "zebra", "giraffe", "backpack", + "umbrella", "handbag", "tie", "suitcase", "frisbee", "skis", + "snowboard", "sports ball", "kite", "baseball bat", "baseball glove", + "skateboard", "surfboard", "tennis racket", "bottle", "wine glass", + "cup", "fork", "knife", "spoon", "bowl", "banana", "apple", "sandwich", + "orange", "broccoli", "carrot", "hot dog", "pizza", "donut", "cake", + "chair", "couch", "potted plant", "bed", "dining table", "toilet", "tv", + "laptop", "mouse", "remote", "keyboard", "cell phone", "microwave", + "oven", "toaster", "sink", "refrigerator", "book", "clock", "vase", + "scissors", "teddy bear", "hair drier", "toothbrush" + }; + coco_class_name = std::vector(coco_class_name_, std::end( coco_class_name_ )); + src = cv::Mat(cv::Size(2,3), CV_32F); + dst = cv::Mat(cv::Size(2,3), CV_32F); + // dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes); + + checkCuda(cudaMallocHost(&input_h, sizeof(dnnType)*netRT->input_dim.tot())); + checkCuda(cudaMallocHost(&input, sizeof(dnnType)*netRT->input_dim.tot())); + checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot())); + + dim_hm = tk::dnn::dataDim_t(1, 80, 56, 56, 1); + dim_wh = tk::dnn::dataDim_t(1, 2, 56, 56, 1); + dim_reg = tk::dnn::dataDim_t(1, 2, 56, 56, 1); + + checkCuda( cudaMalloc(&topk_scores, dim_hm.c * K *sizeof(float)) ); + checkCuda( cudaMalloc(&topk_inds_, dim_hm.c * K *sizeof(int)) ); + checkCuda( cudaMalloc(&topk_ys_, dim_hm.c * K *sizeof(float)) ); + checkCuda( cudaMalloc(&topk_xs_, dim_hm.c * K *sizeof(float)) ); + checkCuda( cudaMalloc(&ids_d, dim_hm.c * dim_hm.h * dim_hm.w*sizeof(int)) ); + checkCuda( cudaMalloc(&ids_2d, dim_hm.c * dim_hm.h * dim_hm.w*sizeof(int)) ); + checkCuda( cudaMallocHost(&ids_, dim_hm.c * dim_hm.h * dim_hm.w*sizeof(int)) ); + checkCuda( cudaMallocHost(&ids_2, dim_hm.c * dim_hm.h * dim_hm.w*sizeof(int)) ); + for(int i =0; iinfer(dim2, input_d); + TIMER_STOP + dim2.print(); + } + // checkResult(dim2.tot(), input_h, input); + std::cout<<" --- pre-process ---\n"; + end_t = std::chrono::steady_clock::now(); + std::cout << " TIME : " << std::chrono::duration_cast(end_t - step_t).count() << " ms" << std::endl; + step_t = end_t; + // ------------------------------------ process -------------------------------------------- + + rt_out[0] = (dnnType *)netRT->buffersRT[1]; + rt_out[1] = (dnnType *)netRT->buffersRT[2]; + rt_out[2] = (dnnType *)netRT->buffersRT[3]; + rt_out[3] = (dnnType *)netRT->buffersRT[4]; + + activationSIGMOIDForward(rt_out[0], rt_out[0], dim_hm.tot()); + checkCuda( cudaDeviceSynchronize() ); + end_t = std::chrono::steady_clock::now(); + std::cout << " TIME sigmoid : " << std::chrono::duration_cast(end_t - step_t).count() << " ms" << std::endl; + step_t = end_t; + + subtractWithThreshold(rt_out[0], rt_out[0] + dim_hm.tot(), rt_out[1], rt_out[0]); + + float *prova; + checkCuda( cudaMallocHost(&prova, K*sizeof(float)) ); + checkCuda( cudaMemcpy(prova, rt_out[0], K*sizeof(float), cudaMemcpyDeviceToHost) ); + std::cout<<"heat:\n"; + for(int i=0; i toll || hm_h[i]-hmax_h[i] < -toll){ + // hm_h[i] = 0.0f; + // } + // } + // checkCuda( cudaFreeHost(hmax_h) ); + std::cout<<" --- hmax ---\n"; + end_t = std::chrono::steady_clock::now(); + std::cout << " TIME : " << std::chrono::duration_cast(end_t - step_t).count() << " ms" << std::endl; + step_t = end_t; + // ----------- nms end + // ----------- topk + + + // thrust::device_vector ids_d; + // int ids[dim_hm.h * dim_hm.w]; + // for(int i=0; i ids2( dim_hm.h * dim_hm.w ); + // for(int i=0; i dim_hm.h * dim_hm.w){ + printf ("Error topk (K is too large)\n"); + return; + } + + + checkCuda( cudaMemcpy(ids_d, ids_, dim_hm.c * dim_hm.h * dim_hm.w*sizeof(int), cudaMemcpyHostToDevice) ); + // checkCuda( cudaMemcpy(ids_2d, ids_2, dim_hm.h * dim_hm.w*sizeof(int), cudaMemcpyHostToDevice) ); + + + // sortAndTopKonDevice(rt_out[0], ids_2d, topk_scores, topk_inds_ , topk_ys_ , topk_xs_ ,dim_hm.h * dim_hm.w, K, dim_hm.c); + // checkCuda( cudaDeviceSynchronize() ); + + // for(int i=0; ioutput_dim.h * hm->output_dim.w elements for each channel and sort it. Then find the first 100 elements + // // memcpy(ids2, ids, dim_hm.h * dim_hm.w); + // sort(rt_out[0]+ i * dim_hm.h * dim_hm.w, + // rt_out[0]+ i * dim_hm.h * dim_hm.w + dim_hm.h * dim_hm.w, + // ids_d); + // // end_t = std::chrono::steady_clock::now(); + // // std::cout << " TIME sort channel "<(end_t - step_t).count() << " ms" << std::endl; + // // step_t = end_t; + // topk(rt_out[0]+ i * dim_hm.h * dim_hm.w, ids_d, K, topk_scores + i*K, + // topk_inds_ + i*K, topk_ys_ + i*K, topk_xs_ + i*K); + // // checkCuda( cudaMemcpy(ids2, ids2_d, dim_hm.h * dim_hm.w*sizeof(int), cudaMemcpyDeviceToHost) ); + + // // for (int j=0; j(end_t - step_t).count() << " ms" << std::endl; + // // step_t = end_t; + + // } + // checkCuda( cudaFree(ids_d )); + std::cout<<" --- a 100 ---\n"; + end_t = std::chrono::steady_clock::now(); + std::cout << " TIME sort topk on 80 channel: " << std::chrono::duration_cast(end_t - step_t).count() << " ms" << std::endl; + step_t = end_t; + // final + + // sort(topk_scores, + // topk_scores + dim_hm.c * K, + // topk_inds_); + sort(rt_out[0], + rt_out[0]+dim_hm.tot(), + ids_d); + checkCuda( cudaDeviceSynchronize() ); + int *topk_inds; + checkCuda( cudaMallocHost(&topk_inds, K*sizeof(int)) ); + // checkCuda( cudaMemcpy(topk_inds, ids_d, K*sizeof(int), cudaMemcpyDeviceToHost) ); + // for(int i=0; i(end_t - step_t).count() << " ms" << std::endl; + step_t = end_t; + + // topk(topk_scores, topk_inds_, K, scores_d, + // topk_inds_d, topk_ys_d, topk_xs_d); + topk(rt_out[0], ids_d, K, scores_d, + topk_inds_d, topk_ys_d, topk_xs_d); + checkCuda( cudaDeviceSynchronize() ); + end_t = std::chrono::steady_clock::now(); + std::cout << " TIME topk channel: " << std::chrono::duration_cast(end_t - step_t).count() << " ms" << std::endl; + step_t = end_t; + + + checkCuda( cudaMemcpy(topk_inds, topk_inds_d, K*sizeof(int), cudaMemcpyDeviceToHost) ); + for(int i=0; i(end_t - step_t).count() << " ms" << std::endl; + step_t = end_t; + // ----------- topk end + + // dnnType *reg_aus; + // checkCuda( cudaMallocHost(®_aus, dim_reg.tot()*sizeof(dnnType)) ); + // checkCuda( cudaMemcpy(reg_aus, rt_out[3], dim_reg.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost) ); + + // for(int i = 0; i < K; i++){ + // topk_xs[i] = topk_xs[i] + reg_aus[topk_inds[i]]; + // topk_ys[i] = topk_ys[i] + reg_aus[topk_inds[i]+dim_reg.h*dim_reg.w]; + // } + topKxyAddOffset(topk_inds_d, K, dim_reg.h*dim_reg.w, inttopk_xs_d, inttopk_ys_d, topk_xs_d, topk_ys_d, rt_out[3]); + // checkCuda( cudaDeviceSynchronize() ); + end_t = std::chrono::steady_clock::now(); + std::cout << " TIME add offset: " << std::chrono::duration_cast(end_t - step_t).count() << " ms" << std::endl; + step_t = end_t; + // checkCuda( cudaFreeHost(reg_aus) ); + + // dnnType *wh_aus; + + // checkCuda( cudaMemcpy(wh_aus, rt_out[2], dim_wh.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost) ); + bboxes(topk_inds_d, K, dim_wh.h*dim_wh.w, topk_xs_d, topk_ys_d, rt_out[2], bbx0_d, bbx1_d, bby0_d, bby1_d); + // checkCuda( cudaDeviceSynchronize() ); + checkCuda( cudaMemcpy(bbx0, bbx0_d, K * sizeof(float), cudaMemcpyDeviceToHost) ); + checkCuda( cudaMemcpy(bby0, bby0_d, K * sizeof(float), cudaMemcpyDeviceToHost) ); + checkCuda( cudaMemcpy(bbx1, bbx1_d, K * sizeof(float), cudaMemcpyDeviceToHost) ); + checkCuda( cudaMemcpy(bby1, bby1_d, K * sizeof(float), cudaMemcpyDeviceToHost) ); + // for(int i = 0; i < K; i++){ + // bboxes[i * 4] = topk_xs[i] - wh_aus[topk_inds[i]] / 2; + // bboxes[i * 4 + 1] = topk_ys[i] - wh_aus[topk_inds[i]+dim_reg.h*dim_reg.w] / 2; + // bboxes[i * 4 + 2] = topk_xs[i] + wh_aus[topk_inds[i]] / 2; + // bboxes[i * 4 + 3] = topk_ys[i] + wh_aus[topk_inds[i]+dim_reg.h*dim_reg.w] / 2; + // } + // for(int i = 0; i < K; i++){ + // std::cout<<"-----\n(x0, y0) = ("<(end_t - step_t).count() << " ms" << std::endl; + step_t = end_t; + // servono [bboxes, scores, clses] + // checkCuda( cudaDeviceSynchronize() ); + + std::cout<<" --- process ---\n"; + end_t = std::chrono::steady_clock::now(); + std::cout << " TIME : " << std::chrono::duration_cast(end_t - step_t).count() << " ms" << std::endl; + step_t = end_t; + // ---------------------------------- post-process ----------------------------------------- + + // --------- ctdet_post_process + // --------- transform_preds + src.at(0,0)=c[0]; + src.at(0,1)=c[1]; + src.at(1,0)=c[0]; + src.at(1,1)=c[1] + s[0] * -0.5; + dst.at(0,0)=width * 0.5; + dst.at(0,1)=width * 0.5; + dst.at(1,0)=width * 0.5; + dst.at(1,1)=width * 0.5 + width * -0.5; + + src.at(2,0)=src.at(1,0) + (-src.at(0,1)+src.at(1,1) ); + src.at(2,1)=src.at(1,1) + (src.at(0,0)-src.at(1,0) ); + dst.at(2,0)=dst.at(1,0) + (-dst.at(0,1)+dst.at(1,1) ); + dst.at(2,1)=dst.at(1,1) + (dst.at(0,0)-dst.at(1,0) ); + + + cv::Mat trans2(cv::Size(3,2), CV_32F); + trans2 = cv::getAffineTransform( dst, src ); + end_t = std::chrono::steady_clock::now(); + std::cout << " TIME getAffineTrans 2: " << std::chrono::duration_cast(end_t - step_t).count() << " ms" << std::endl; + step_t = end_t; + cv::Mat new_pt1(cv::Size(1,2), CV_32F); + cv::Mat new_pt2(cv::Size(1,2), CV_32F); + + + for(int i = 0; i(0,0)=static_cast(trans2.at(0,0))*bbx0[i] + + static_cast(trans2.at(0,1))*bby0[i] + + static_cast(trans2.at(0,2))*1.0; + new_pt1.at(0,1)=static_cast(trans2.at(1,0))*bbx0[i] + + static_cast(trans2.at(1,1))*bby0[i] + + static_cast(trans2.at(1,2))*1.0; + + new_pt2.at(0,0)=static_cast(trans2.at(0,0))*bbx1[i] + + static_cast(trans2.at(0,1))*bby1[i] + + static_cast(trans2.at(0,2))*1.0; + new_pt2.at(0,1)=static_cast(trans2.at(1,0))*bbx1[i] + + static_cast(trans2.at(1,1))*bby1[i] + + static_cast(trans2.at(1,2))*1.0; + + // std::cout<<"\n new: "<(0,0); + target_coords[i*4+1] = new_pt1.at(0,1); + target_coords[i*4+2] = new_pt2.at(0,0); + target_coords[i*4+3] = new_pt2.at(0,1); + // std::cout<(0,0)<<", "<(0,1)<<", "<(0,0)<<", "<(0,1)< thresh){ + std::cout<<"th: "<(end_t - step_t).count() << " ms" << std::endl; + step_t = end_t; + std::cout<<"TOTAL: \n"; + TIMER_STOP +} +}} \ No newline at end of file diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 31d57fc..33e9d62 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -292,7 +292,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) { if(l->pool_mode == tkdnnPoolingMode_t::POOLING_AVERAGE_EXCLUDE_PADDING) ptype = PoolingType::kMAX_AVERAGE_BLEND; - if(l->input_dim.h % 2 == 1 && l->input_dim.w % 2 == 1) + // if(l->input_dim.h % 2 == 1 && l->input_dim.w % 2 == 1) + if(l->input_dim.h == l->output_dim.h && l->input_dim.w == l->output_dim.w) { IPlugin *plugin = new ResizeLayerRT( l->output_dim.c,l->output_dim.h+1,l->output_dim.w+1 ); IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); @@ -324,6 +325,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kRELU); checkNULL(lRT); return lRT; + } else if(l->act_mode == CUDNN_ACTIVATION_SIGMOID) { + IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kSIGMOID); + checkNULL(lRT); + return lRT; } else { FatalError("this Activation mode is not yet implemented"); diff --git a/src/Pooling.cpp b/src/Pooling.cpp index 236adbf..1ad0664 100644 --- a/src/Pooling.cpp +++ b/src/Pooling.cpp @@ -7,8 +7,8 @@ namespace tk { namespace dnn { Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW, int paddingH, int paddingW, - tkdnnPoolingMode_t pool_mode) : - Layer(net) { + tkdnnPoolingMode_t pool_mode, bool final) : + Layer(net, final) { this->winH = winH; this->winW = winW; @@ -53,10 +53,17 @@ Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW, //compute w and h as in darknet int padH = paddingH == 0? winH -1 : paddingH; int padW = paddingW == 0? winW -1 : paddingW; - - h = (h + padH - winH)/strideH +1; - w = (w + padW - winW)/strideW +1; - + + if(final){ + h = (h + padH - winH)/strideH +1 +1; + w = (w + padW - winW)/strideW +1 +1; + } + else{ + h = (h + padH - winH)/strideH +1; + w = (w + padW - winW)/strideW +1; + } + + // h = (h + winH*this->paddingH)/strideH; // w = (w + winW*this->paddingW)/strideW; diff --git a/src/sorting.cu b/src/sorting.cu new file mode 100644 index 0000000..e007be1 --- /dev/null +++ b/src/sorting.cu @@ -0,0 +1,104 @@ + +#include "sorting.h" + +void sort(dnnType *src_begin, dnnType *src_end, int *idsrc) +{ + thrust::sort_by_key(thrust::device, + src_begin, src_end, idsrc, + thrust::greater()); + + // thrust::stable_sort_by_key(thrust::device, + // src_begin, src_end, idsrc, + // thrust::greater()); +} + +void topk(dnnType *src_begin, int *idsrc, int K, float *topk_scores, + int *topk_inds, float *topk_ys, float *topk_xs) +{ + checkCuda( cudaMemcpy(topk_scores, (float *)src_begin, K*sizeof(float), cudaMemcpyDeviceToDevice) ); + checkCuda( cudaMemcpy(topk_inds, idsrc, K*sizeof(int), cudaMemcpyDeviceToDevice) ); + // topk_ys_[i*K +count] = (int)(ids2[j] / width); + // topk_xs_[i*K +count] = (int)(ids2[j] % width); + +} + +__global__ +void sortAndTopK_kernel(dnnType *src_begin, int *idsrc, float *topk_scores, int *topk_inds, float *topk_ys, float *topk_xs,const int size, const int K){ + int i = blockDim.x*blockIdx.x + threadIdx.x; + + thrust::sort_by_key(thrust::device, src_begin + i * size, src_begin + i * size + size, idsrc + i * size, thrust::greater()); + thrust::copy_n(thrust::device, src_begin + i * size, K, topk_scores + i * K); + // thrust::copy_n(thrust::device, idsrc + i * size, K, topk_inds + i * K ); + thrust::copy_n(thrust::device, idsrc + i * size, K, topk_inds + i * K ); +} + +void sortAndTopKonDevice(dnnType *src_begin, int *idsrc, float *topk_scores, int *topk_inds, float *topk_ys, float *topk_xs, const int size, const int K, const int n_classes) +{ + int blocks = n_classes; + int threads = 1; + + sortAndTopK_kernel<<>>(src_begin, idsrc, topk_scores, topk_inds, topk_ys, topk_xs, size, K); + +} + +struct threshold : public thrust::binary_function +{ + __host__ __device__ + float operator()(float x, float y) { + float toll = 1e-6; + if(fabsf(x-y)>toll) + return 0.0f; + else + return x; + } +}; + +void subtractWithThreshold(dnnType *src_begin, dnnType *src_end, dnnType *src2_begin, dnnType *src_out){ + struct threshold op; + thrust::transform(thrust::device, src_begin, src_end, src2_begin, src_out, op); +} + +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()); + thrust::transform(thrust::device, ids_begin, ids_end, thrust::make_constant_iterator(size), ys, thrust::divides()); + thrust::transform(thrust::device, ids_begin, ids_end, thrust::make_constant_iterator(size), xs, thrust::modulus()); + +} + +void topKxyAddOffset(int * ids_begin, const int K, const int size, int *intxs_begin, int *intys_begin, float *xs_begin, float *ys_begin, dnnType *src_begin){ + float *src_out; + checkCuda( cudaMalloc(&src_out, K *sizeof(float)) ); + thrust::gather(thrust::device, ids_begin, ids_begin + K, src_begin, src_out); + thrust::transform(thrust::device, intxs_begin, intxs_begin + K, src_out, xs_begin, thrust::plus()); + int *ids_out; + checkCuda( cudaMalloc(&ids_out, K *sizeof(int)) ); + thrust::transform(thrust::device, ids_begin, ids_begin + K, thrust::make_constant_iterator(size), ids_out, thrust::plus()); + thrust::gather(thrust::device, ids_out, ids_out+K, src_begin, src_out); + thrust::transform(thrust::device, intys_begin, intys_begin + K, src_out, ys_begin, thrust::plus()); + checkCuda( cudaFree(src_out) ); + checkCuda( cudaFree(ids_out) ); +} + +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; + checkCuda( cudaMalloc(&src_out, K *sizeof(float)) ); + thrust::gather(thrust::device, ids_begin, ids_begin + K, src_begin, src_out); + thrust::transform(thrust::device, src_out, src_out + K, thrust::make_constant_iterator(2), src_out, thrust::divides()); + // x0 + thrust::transform(thrust::device, xs_begin, xs_begin + K, src_out, bbx0, thrust::minus()); + // x1 + thrust::transform(thrust::device, xs_begin, xs_begin + K, src_out, bbx1, thrust::plus()); + int *ids_out; + checkCuda( cudaMalloc(&ids_out, K *sizeof(int)) ); + thrust::transform(thrust::device, ids_begin, ids_begin + K, thrust::make_constant_iterator(size), ids_out, thrust::plus()); + thrust::gather(thrust::device, ids_out, ids_out + K, src_begin, src_out); + thrust::transform(thrust::device, src_out, src_out + K, thrust::make_constant_iterator(2), src_out, thrust::divides()); + // y0 + thrust::transform(thrust::device, ys_begin, ys_begin + K, src_out, bby0, thrust::minus()); + // y1 + thrust::transform(thrust::device, ys_begin, ys_begin + K, src_out, bby1, thrust::plus()); + checkCuda( cudaFree(src_out) ); + checkCuda( cudaFree(ids_out) ); +} + diff --git a/tests/resnet101_cnet/resnet101_cnet.cpp b/tests/resnet101_cnet/resnet101_cnet.cpp index c89591a..061e4b6 100644 --- a/tests/resnet101_cnet/resnet101_cnet.cpp +++ b/tests/resnet101_cnet/resnet101_cnet.cpp @@ -1,5 +1,12 @@ #include + +#include "kernels.h" +#include "Yolo3Detection.h" #include "tkdnn.h" +#include +#include // std::iota +#include // std::sort +// #include "utils.h" const char *input_bin = "../tests/resnet101_cnet/debug/input.bin"; const char *conv1_bin = "../tests/resnet101_cnet/layers/conv1.bin"; @@ -172,9 +179,56 @@ const char *reg_conv2_bin = "../tests/resnet101_cnet/layers/reg-2.bin"; const char *fc_bin = "../tests/resnet101_cnet/layers/fc.bin"; const char *output_bin[]={ - "../tests/resnet101_cnet/debug/hm.bin", - "../tests/resnet101_cnet/debug/wh.bin", - "../tests/resnet101_cnet/debug/reg.bin"}; +"../tests/resnet101_cnet/debug/hm.bin", +"../tests/resnet101_cnet/debug/wh.bin", +"../tests/resnet101_cnet/debug/reg.bin"}; + + + +std::vector sort_indexes(const std::vector &v) { + + // initialize original index locations + std::vector idx(v.size()); + iota(idx.begin(), idx.end(), 0); + + // sort indexes based on comparing values in v + sort(idx.begin(), idx.end(), + [&v](size_t i1, size_t i2) {return v[i1] > v[i2];}); + + return idx; +} + +float _colors[6][3] = { {1,0,1}, {0,0,1},{0,1,1},{0,1,0},{1,1,0},{1,0,0} }; +float get_color(int c, int x, int max) +{ + float ratio = ((float)x/max)*5; + int i = floor(ratio); + int j = ceil(ratio); + ratio -= i; + float r = (1-ratio) * _colors[i % 6][c % 3] + ratio*_colors[j % 6][c % 3]; + //printf("%f\n", r); + return r; +} + +int computeDetections(dnnType *hm_d, dnnType *wh_d, dnnType *reg_d, int hm_dim, int wh_dim, int reg_dim, bool cat_spec_wh, int k){ + // _nms + int kernel = 3; + int pad = (kernel - 1)/2; + std::cout<<"computeDetections\n"; + // dnnType *hmax; + // tk::dnn::Pooling maxpool(&hmax, 3, 3, 2, 2, 1, 1, tk::dnn::POOLING_MAX) + // = (dnnType *) + + // net.functional.max_pool2d( + // heat, (kernel, kernel), stride=1, padding=pad) + // keep = (hmax == heat).float() + // return heat * keep +} + +int process(dnnType *hm_d, dnnType *wh_d, dnnType *reg_d, int hm_dim, int wh_dim, int reg_dim){ + std::cout<<"process\n"; + // computeDetections(hm_d, wh_d, reg_d, hm_dim, wh_dim, reg_dim, false, 100); +} int main() { @@ -211,7 +265,6 @@ int main() last = layer1_0_relu; } - // tk::dnn::Activation *last_activation = (tk::dnn::Activation *) net.layers[net.num_layers-1]; // layer 2 int id_layer2_bin = 0; for(int i=0; i<4;i++) @@ -318,6 +371,10 @@ int main() tk::dnn::Conv2d *hm_conv1 = new tk::dnn::Conv2d(&net, 64, 3, 3, 1, 1, 1, 1, hm_conv1_bin, false); tk::dnn::Activation *hm_relu1 = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); tk::dnn::Conv2d *hm = new tk::dnn::Conv2d(&net, 80, 1, 1, 1, 1, 0, 0, hm_conv2_bin, false, false, true); + int kernel = 3; + int pad = (kernel - 1)/2; + tk::dnn::Activation *hm_sig = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_SIGMOID); + tk::dnn::Pooling *hmax = new tk::dnn::Pooling(&net, kernel, kernel, 1, 1, pad, pad, tk::dnn::POOLING_MAX, true); tk::dnn::Route *route_1_0 = new tk::dnn::Route(&net, route_1_0_layers, 1); tk::dnn::Conv2d *wh_conv1 = new tk::dnn::Conv2d(&net, 64, 3, 3, 1, 1, 1, 1, wh_conv1_bin, false); @@ -365,7 +422,7 @@ int main() } tk::dnn::Layer *outs[3] = { hm, wh, reg }; - + int out_count = 1; for(int i=0; i<3; i++) { printCenteredTitle((std::string(" RESNET CHECK RESULTS ") + std::to_string(i) + " ").c_str(), '=', 30); @@ -381,7 +438,10 @@ int main() dnnType *cudnn_out, *rt_out; cudnn_out = outs[i]->dstData; - rt_out = (dnnType *)netRT.buffersRT[i+1]; + rt_out = (dnnType *)netRT.buffersRT[i+out_count]; + // there is the maxpool. It isn't an output but it is necessary for the process section + if(i==0) + out_count ++; std::cout << "CUDNN vs correct"; checkResult(odim, cudnn_out, out); @@ -391,5 +451,636 @@ int main() std::cout << "CUDNN vs TRT "; checkResult(odim, cudnn_out, rt_out); } + + TIMER_START + + // -------- transofrm compose + cv::Mat imageOrig = cv::imread("/media/davide/DATA/shared_home/Projects/Professionale/repos/photo_2020-01-14_09-56-07.jpg"); + cv::Mat imageF; + imageOrig.convertTo(imageF, CV_32FC3, 1/255.0); + cv::Mat image; + cv::Size sz = imageF.size(); + std::cout<<"image: "<output_dim.tot()*sizeof(dnnType)) ); + + dnnType *rt_out[4]; + rt_out[0] = (dnnType *)netRT.buffersRT[1]; + rt_out[1] = (dnnType *)netRT.buffersRT[2]; + rt_out[2] = (dnnType *)netRT.buffersRT[3]; + rt_out[3] = (dnnType *)netRT.buffersRT[4]; + + // checkCuda( cudaMemcpy(hm_h, rt_out[0], hm->output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost) ); + std::cout<<"hm\n"; + hm->output_dim.print(); + // for(int i=0; ioutput_dim.tot(); i++ ){ + // std::cout<output_dim.tot()); + checkCuda( cudaDeviceSynchronize() ); + + + checkCuda( cudaMemcpy(hm_h, rt_out[0], hm->output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost) ); + std::cout<<"hm\n"; + hm->output_dim.print(); + // for(int i=0; ioutput_dim.tot(); i++ ){ + // std::cout<infer(hmax->input_dim.tot(), rt_out[0]); + // keep = (hmax == heat).float() + // return heat * keep + + dnnType *hmax_h; + checkCuda( cudaMallocHost(&hmax_h, hmax->output_dim.tot()*sizeof(dnnType)) ); + checkCuda( cudaMemcpy(hmax_h, rt_out[1], hmax->output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost) ); + + std::cout<<"hmax\n"; + hmax->output_dim.print(); + // for(int i=0; ioutput_dim.tot(); i++ ){ + // std::cout<output_dim.print(); + std::cout<<"hmax:\n"; + hmax->output_dim.print(); + // return 0; + float toll = 0.000001; + for(int i=0; i < hm->output_dim.tot(); i++){ + if(hm_h[i]-hmax_h[i] > toll || hm_h[i]-hmax_h[i] < -toll){ + hm_h[i] = 0.0f; + } + } + // std::cout<<"\n"; + // for(int i=0; ioutput_dim.tot(); i++ ){ + // std::cout<dstData, hm_h, hm->output_dim.tot()*sizeof(dnnType), cudaMemcpyHostToDevice) ); + checkCuda( cudaFreeHost(hmax_h) ); + // ----------- nms end + // ----------- topk + int K = 100; + int width = 56; // TODO + float *topk_scores; + int *topk_inds_; + float *topk_ys_; + float *topk_xs_; + std::cout<<"mah: "<output_dim.c * K<output_dim.c * K *sizeof(float)) ); + checkCuda( cudaMallocHost(&topk_inds_, hm->output_dim.c * K *sizeof(int)) ); + checkCuda( cudaMallocHost(&topk_ys_, hm->output_dim.c * K *sizeof(float)) ); + checkCuda( cudaMallocHost(&topk_xs_, hm->output_dim.c * K *sizeof(float)) ); + std::cout<<"1\n"; + dnnType *hm_aus; + checkCuda( cudaMallocHost(&hm_aus, hm->output_dim.h * hm->output_dim.w *sizeof(dnnType)) ); + std::cout<<"2\n"; + int count; + std::vector v = {2.0, 3.0, 9.0}; + for (auto i: sort_indexes(v)) { + std::cout << i<< "--" <output_dim.c; i++){ + count = 0; + // get the hm->output_dim.h * hm->output_dim.w elements for each channel and sort it. Then find the first 100 elements + checkCuda( cudaMemcpy(hm_aus, hm_h + i * hm->output_dim.h * hm->output_dim.w, + hm->output_dim.h * hm->output_dim.w * sizeof(dnnType), cudaMemcpyHostToHost) ); + // std::cout<<"top scores: "<output_dim.h * hm->output_dim.w<<"\n"; + // for(int k=0; koutput_dim.h * hm->output_dim.w; k++) + // std::cout< my_vector {arr, arr + arr_length} + std::vector my_vector{hm_aus, hm_aus + hm->output_dim.h * hm->output_dim.w}; + for (auto j: sort_indexes(my_vector)) { + // std::cout <<"j: "< "<< hm_aus[j] << std::endl; + topk_scores[i*K + count] = hm_aus[j]; + topk_inds_[i*K +count] = j; + topk_ys_[i*K +count] = (int)(j / width); + topk_xs_[i*K +count] = (int)(j % width); + if(++count == K) + break; + } + } + std::cout<<"topk_xs_[0]: "<output_dim.c * K; i++) + std::cout< my_vector{topk_scores, topk_scores + hm->output_dim.c * K }; + for (auto j: sort_indexes(my_vector)) { + // std::cout <<"j: "< "<< hm_aus[j] << std::endl; + scores[count] = topk_scores[j]; + clses[count] = (int)(j / K); + topk_inds[count] = topk_inds_[j]; + topk_ys[count] = topk_ys_[j]; + topk_xs[count] = topk_xs_[j]; + if(++count == K) + break; + } + checkCuda( cudaFreeHost(topk_scores) ); + checkCuda( cudaFreeHost(topk_inds_) ); + checkCuda( cudaFreeHost(topk_ys_) ); + checkCuda( cudaFreeHost(topk_xs_) ); + std::cout<<"5\n"; + // ----------- topk end + std::cout<<"topk_xs[0]: "<output_dim.tot()*sizeof(dnnType)) ); + checkCuda( cudaMemcpy(reg_aus, rt_out[3], reg->output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost) ); + std::cout<<"reg:\n"; + reg->output_dim.print(); + // for(int i=0; ioutput_dim.tot(); i++ ){ + // std::cout<output_dim.h*reg->output_dim.w]; + } + std::cout<<"topk_xs[0]: "<output_dim.tot()*sizeof(dnnType)) ); + checkCuda( cudaMallocHost(&bboxes, 4 * K *sizeof(dnnType)) ); + checkCuda( cudaMemcpy(wh_aus, rt_out[2], wh->output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost) ); + std::cout<<"7\n"; + for(int i = 0; i< K; i++) + std::cout<output_dim.h*reg->output_dim.w] / 2; + bboxes[i * 4 + 2] = topk_xs[i] + wh_aus[topk_inds[i]] / 2; + bboxes[i * 4 + 3] = topk_ys[i] + wh_aus[topk_inds[i]+reg->output_dim.h*reg->output_dim.w] / 2; + } + ////////////////// fin qui ok + + checkCuda( cudaFreeHost(wh_aus) ); + checkCuda( cudaFreeHost(topk_inds) ); + checkCuda( cudaFreeHost(topk_ys) ); + checkCuda( cudaFreeHost(topk_xs) ); + + std::cout<<"8\n"; + float *detections; + std::cout<<"bboxes:\n"; + for(int i = 0; i < K+1; i++){ + std::cout<(0,0)=c[0]; + src.at(0,1)=c[1]; + src.at(1,0)=c[0]; + src.at(1,1)=c[1] + s[0] * -0.5; + dst.at(0,0)=width * 0.5; + dst.at(0,1)=width * 0.5; + dst.at(1,0)=width * 0.5; + dst.at(1,1)=width * 0.5 + width * -0.5; + + src.at(2,0)=src.at(1,0) + (-src.at(0,1)+src.at(1,1) ); + src.at(2,1)=src.at(1,1) + (src.at(0,0)-src.at(1,0) ); + dst.at(2,0)=dst.at(1,0) + (-dst.at(0,1)+dst.at(1,1) ); + dst.at(2,1)=dst.at(1,1) + (dst.at(0,0)-dst.at(1,0) ); + std::cout<<"src: "<(0,0)<<" - "<(0,1)<<" - "<(0,2)<<"\n"<(1,0)<<" - "<(1,1)<<" - "<(1,2)<(0,0)=detections[i*4]; + // new_pt1.at(0,1)=detections[i*4+1]; + // new_pt1.at(0,2)=1.0; + // new_pt1 << detections[i], detections[i+K], 1.0; + // std::cout<<"----\ni: "<(0,0)=static_cast(trans2.at(0,0))*detections[i*4] + + static_cast(trans2.at(0,1))*detections[i*4+1] + + static_cast(trans2.at(0,2))*1.0; + new_pt1.at(0,1)=static_cast(trans2.at(1,0))*detections[i*4] + + static_cast(trans2.at(1,1))*detections[i*4+1] + + static_cast(trans2.at(1,2))*1.0; + + new_pt2.at(0,0)=static_cast(trans2.at(0,0))*detections[i*4+2] + + static_cast(trans2.at(0,1))*detections[i*4+3] + + static_cast(trans2.at(0,2))*1.0; + new_pt2.at(0,1)=static_cast(trans2.at(1,0))*detections[i*4+2] + + static_cast(trans2.at(1,1))*detections[i*4+3] + + static_cast(trans2.at(1,2))*1.0; + + + // std::cout<<"\n new: "<(0,0); + target_coords[i*4+1] = new_pt1.at(0,1); + target_coords[i*4+2] = new_pt2.at(0,0); + target_coords[i*4+3] = new_pt2.at(0,1); + // std::cout<(0,0)<<", "<(0,1)<<", "<(0,0)<<", "<(0,1)< coco_class_name(coco_class_name_, std::end( coco_class_name_ )); + int num_classes = 80; + float vis_threshold = 0.3; + // int *classes; + std::vector detected; + // checkCuda( cudaMallocHost(&classes, K *sizeof(int)) ); + // checkCuda( cudaMemcpy(classes, detections + 5 * K *sizeof(dnnType), K *sizeof(dnnType), cudaMemcpyHostToHost) ); + for(int i = 0; i i+1 (1:80); + + if(scores[j] > vis_threshold){ + std::cout<<"th: "<