From 266330009c42479e1e4f32a064b7ea6ededf02c5 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 10 Aug 2017 16:22:17 +0200 Subject: [PATCH] opencv viz --- CMakeLists.txt | 11 ++++- include/Layer.h | 47 +++++++++++---------- include/Network.h | 2 +- include/NetworkRT.h | 4 +- include/kernels.h | 8 ++-- include/utils.h | 16 ++++---- src/Activation.cpp | 8 ++-- src/Conv2d.cpp | 12 +++--- src/Dense.cpp | 8 ++-- src/Flatten.cpp | 4 +- src/MulAdd.cpp | 12 +++--- src/Network.cpp | 2 +- src/NetworkRT.cpp | 10 ++--- src/Pooling.cpp | 16 ++++---- src/Region.cpp | 63 +++++++++++++++++++++++++---- src/Reorg.cpp | 4 +- src/Route.cpp | 8 ++-- src/Softmax.cpp | 8 ++-- src/kernels/activation_elu.cu | 6 +-- src/kernels/activation_leaky.cu | 4 +- src/kernels/activation_logistic.cu | 4 +- src/kernels/reorg.cu | 2 +- src/pluginsRT/ActivationLeakyRT.cpp | 4 +- src/pluginsRT/RegionRT.cpp | 6 +-- src/pluginsRT/ReorgRT.cpp | 4 +- src/utils.cpp | 46 ++++++++++----------- tests/mnist/test_mnist.cpp | 10 ++--- tests/mnist/test_mnistRT.cpp | 8 ++-- tests/simple/test_simple.cpp | 8 ++-- tests/yolo-tiny/yolo-tiny.cpp | 9 +++-- tests/yolo/yolo.cpp | 9 +++-- 31 files changed, 211 insertions(+), 152 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 87da1be..afc523c 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -18,6 +18,13 @@ if(DEBUG) endif() find_package(CUDA QUIET REQUIRED) +find_package(OpenCV QUIET) +if(${OpenCV_FOUND}) + message("Compiling with openCV support") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV") +else() + message(WARNING "OpenCV not found, compiling without it") +endif() cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS}) cuda_add_library(kernels SHARED src/kernels/activation_elu.cu @@ -27,11 +34,11 @@ cuda_add_library(kernels SHARED src/kernels/activation_elu.cu src/kernels/softmax.cu) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -std=c++11") -include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS}) +include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS}) add_library(tkDNN SHARED src/Layer.cpp src/LayerWgs.cpp src/Dense.cpp src/Activation.cpp src/Conv2d.cpp src/Flatten.cpp src/MulAdd.cpp src/Pooling.cpp src/Softmax.cpp src/Route.cpp src/Reorg.cpp src/Region.cpp src/Network.cpp src/utils.cpp src/NetworkRT.cpp) -target_link_libraries(tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn -lnvinfer) +target_link_libraries(tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn -lnvinfer ${OpenCV_LIBS}) add_executable(test_simple tests/simple/test_simple.cpp) target_link_libraries(test_simple tkDNN) diff --git a/include/Layer.h b/include/Layer.h index 7496236..aa7d7ca 100644 --- a/include/Layer.h +++ b/include/Layer.h @@ -30,13 +30,13 @@ public: virtual ~Layer(); virtual layerType_t getLayerType() = 0; - virtual value_type* infer(dataDim_t &dim, value_type* srcData) { + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData) { std::cout<<"No infer action for this layer\n"; return NULL; } dataDim_t input_dim, output_dim; - value_type *dstData; //where results will be putted + dnnType *dstData; //where results will be putted std::string getLayerName() { layerType_t type = getLayerType(); @@ -75,14 +75,14 @@ public: int inputs, outputs; std::string weights_path; - value_type *data_h, *data_d; - value_type *bias_h, *bias_d; + dnnType *data_h, *data_d; + dnnType *bias_h, *bias_d; //batchnorm bool batchnorm; - value_type *scales_h, *scales_d; - value_type *mean_h, *mean_d; - value_type *variance_h, *variance_d; + dnnType *scales_h, *scales_d; + dnnType *mean_h, *mean_d; + dnnType *variance_h, *variance_d; }; @@ -96,7 +96,7 @@ public: virtual ~Dense(); virtual layerType_t getLayerType() { return LAYER_DENSE; }; - virtual value_type* infer(dataDim_t &dim, value_type* srcData); + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); }; @@ -120,7 +120,7 @@ public: virtual ~Activation(); virtual layerType_t getLayerType() { return LAYER_ACTIVATION; }; - virtual value_type* infer(dataDim_t &dim, value_type* srcData); + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); protected: cudnnActivationDescriptor_t activDesc; @@ -139,7 +139,7 @@ public: virtual ~Conv2d(); virtual layerType_t getLayerType() { return LAYER_CONV2D; }; - virtual value_type* infer(dataDim_t &dim, value_type* srcData); + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); int kernelH, kernelW, strideH, strideW, paddingH, paddingW; @@ -165,7 +165,7 @@ public: virtual ~Flatten(); virtual layerType_t getLayerType() { return LAYER_FLATTEN; }; - virtual value_type* infer(dataDim_t &dim, value_type* srcData); + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); }; @@ -176,15 +176,15 @@ public: class MulAdd : public Layer { public: - MulAdd(Network *net, value_type mul, value_type add); + MulAdd(Network *net, dnnType mul, dnnType add); virtual ~MulAdd(); virtual layerType_t getLayerType() { return LAYER_MULADD; }; - virtual value_type* infer(dataDim_t &dim, value_type* srcData); + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); protected: - value_type mul, add; - value_type *add_vector; + dnnType mul, add; + dnnType *add_vector; }; @@ -214,13 +214,13 @@ public: virtual ~Pooling(); virtual layerType_t getLayerType() { return LAYER_POOLING; }; - virtual value_type* infer(dataDim_t &dim, value_type* srcData); + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); protected: cudnnPoolingDescriptor_t poolingDesc; tkdnnPoolingMode_t pool_mode; - value_type *tmpInputData, *tmpOutputData; + dnnType *tmpInputData, *tmpOutputData; bool poolOn3d; }; @@ -234,7 +234,7 @@ public: virtual ~Softmax(); virtual layerType_t getLayerType() { return LAYER_SOFTMAX; }; - virtual value_type* infer(dataDim_t &dim, value_type* srcData); + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); }; /** @@ -248,7 +248,7 @@ public: virtual ~Route(); virtual layerType_t getLayerType() { return LAYER_ROUTE; }; - virtual value_type* infer(dataDim_t &dim, value_type* srcData); + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); public: Layer **layers; //ids of layers to be merged @@ -267,7 +267,7 @@ public: virtual ~Reorg(); virtual layerType_t getLayerType() { return LAYER_REORG; }; - virtual value_type* infer(dataDim_t &dim, value_type* srcData); + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); int stride; }; @@ -288,11 +288,13 @@ public: virtual ~Region(); virtual layerType_t getLayerType() { return LAYER_REGION; }; - virtual value_type* infer(dataDim_t &dim, value_type* srcData); + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); - value_type *bias_h, *bias_d; + dnnType *bias_h, *bias_d; int classes, coords, num; float thresh; + 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); @@ -301,6 +303,7 @@ public: 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 showImageResult(dnnType *input_h); }; diff --git a/include/Network.h b/include/Network.h index e0decdb..d948515 100644 --- a/include/Network.h +++ b/include/Network.h @@ -43,7 +43,7 @@ public: /** Do inferece for every added layer */ - value_type* infer(dataDim_t &dim, value_type* data); + dnnType* infer(dataDim_t &dim, dnnType* data); bool addLayer(Layer *l); void print(); diff --git a/include/NetworkRT.h b/include/NetworkRT.h index 98bed30..0fc9a50 100644 --- a/include/NetworkRT.h +++ b/include/NetworkRT.h @@ -21,7 +21,7 @@ public: int buf_input_idx, buf_output_idx; dataDim_t output_dim; - value_type *output; + dnnType *output; cudaStream_t stream; NetworkRT(Network *net); @@ -30,7 +30,7 @@ public: /** Do inferece */ - value_type* infer(dataDim_t &dim, value_type* data); + dnnType* infer(dataDim_t &dim, dnnType* data); nvinfer1::ITensor* convert_layer(nvinfer1::ITensor *input, Layer *l); nvinfer1::ITensor* convert_layer(nvinfer1::ITensor *input, Conv2d *l); diff --git a/include/kernels.h b/include/kernels.h index 97f81c6..088e235 100644 --- a/include/kernels.h +++ b/include/kernels.h @@ -1,10 +1,10 @@ #include "utils.h" -void activationELUForward(value_type* srcData, value_type* dstData, int size); -void activationLEAKYForward(value_type* srcData, value_type* dstData, int size); -void activationLOGISTICForward(value_type* srcData, value_type* dstData, int size); +void activationELUForward(dnnType* srcData, dnnType* dstData, int size); +void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size); +void activationLOGISTICForward(dnnType* srcData, dnnType* dstData, int size); -void reorgForward( value_type* srcData, value_type* dstData, +void reorgForward( dnnType* srcData, dnnType* dstData, int n, int c, int h, int w, int stride); void softmaxForward(float *input, int n, int batch, int batch_offset, int groups, int group_offset, int stride, float temp, float *output); diff --git a/include/utils.h b/include/utils.h index 2eb1c10..caf8def 100644 --- a/include/utils.h +++ b/include/utils.h @@ -12,7 +12,7 @@ #include #include -#define value_type float +#define dnnType float // Colored output #define COL_END "\033[0m" @@ -88,13 +88,13 @@ } void printCenteredTitle(const char *title, char fill, int dim); -void readBinaryFile(const char* fname, int size, value_type** data_h, value_type** data_d, int seek = 0); -int checkResult(int size, value_type *data_d, value_type *correct_d, bool device = true); -void printDeviceVector(int size, value_type* vec_d, bool device = true); -void resize(int size, value_type **data); +void readBinaryFile(const char* fname, int size, dnnType** data_h, dnnType** data_d, int seek = 0); +int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device = true); +void printDeviceVector(int size, dnnType* vec_d, bool device = true); +void resize(int size, dnnType **data); -void matrixTranspose(cublasHandle_t handle, value_type* srcData, value_type* dstData, int rows, int cols); +void matrixTranspose(cublasHandle_t handle, dnnType* srcData, dnnType* dstData, int rows, int cols); -void matrixMulAdd( cublasHandle_t handle, value_type* srcData, value_type* dstData, - value_type* add_vector, int dim, value_type mul); +void matrixMulAdd( cublasHandle_t handle, dnnType* srcData, dnnType* dstData, + dnnType* add_vector, int dim, dnnType mul); #endif //UTILS_H diff --git a/src/Activation.cpp b/src/Activation.cpp index b3716fa..cc91a82 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -9,7 +9,7 @@ Activation::Activation(Network *net, int act_mode) : Layer(net) { this->act_mode = act_mode; - checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) ); + checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) ); if(int(act_mode) < 100) { @@ -43,14 +43,14 @@ Activation::~Activation() { checkCUDNN( cudnnDestroyActivationDescriptor(activDesc) ); } -value_type* Activation::infer(dataDim_t &dim, value_type* srcData) { +dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) { if(act_mode == ACTIVATION_LEAKY) { activationLEAKYForward(srcData, dstData, dim.tot()); } else { - value_type alpha = value_type(1); - value_type beta = value_type(0); + dnnType alpha = dnnType(1); + dnnType beta = dnnType(0); checkCUDNN( cudnnActivationForward(net->cudnnHandle, activDesc, &alpha, diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index e553baf..f1d9a5d 100644 --- a/src/Conv2d.cpp +++ b/src/Conv2d.cpp @@ -76,7 +76,7 @@ Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW, output_dim.l = 1; //allocate data for infer result - checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) ); + checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); } Conv2d::~Conv2d() { @@ -91,12 +91,12 @@ Conv2d::~Conv2d() { checkCuda( cudaFree(dstData) ); } -value_type* Conv2d::infer(dataDim_t &dim, value_type* srcData) { +dnnType* Conv2d::infer(dataDim_t &dim, dnnType* srcData) { // convolution - value_type alpha = value_type(1); - value_type beta = value_type(0); + dnnType alpha = dnnType(1); + dnnType beta = dnnType(0); checkCUDNN( cudnnConvolutionForward(net->cudnnHandle, &alpha, srcTensorDesc, srcData, filterDesc, data_d, convDesc, algo, workSpace, ws_sizeInBytes, @@ -104,8 +104,8 @@ value_type* Conv2d::infer(dataDim_t &dim, value_type* srcData) { if(!batchnorm) { // bias - alpha = value_type(1); - beta = value_type(1); + alpha = dnnType(1); + beta = dnnType(1); checkCUDNN( cudnnAddTensor(net->cudnnHandle, &alpha, biasTensorDesc, bias_d, &beta, dstTensorDesc, dstData) ); diff --git a/src/Dense.cpp b/src/Dense.cpp index fa49a28..237dfc8 100644 --- a/src/Dense.cpp +++ b/src/Dense.cpp @@ -14,7 +14,7 @@ Dense::Dense(Network *net, int out_ch, const char* fname_weights) : output_dim.l = 1; //allocate data for infer result - checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) ); + checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); } Dense::~Dense() { @@ -22,7 +22,7 @@ Dense::~Dense() { checkCuda( cudaFree(dstData) ); } -value_type* Dense::infer(dataDim_t &dim, value_type* srcData) { +dnnType* Dense::infer(dataDim_t &dim, dnnType* srcData) { if (dim.n != 1) FatalError("Not Implemented"); @@ -33,9 +33,9 @@ value_type* Dense::infer(dataDim_t &dim, value_type* srcData) { if (dim_x != input_dim.tot()) FatalError("Input mismatch"); - value_type alpha = value_type(1), beta = value_type(1); + dnnType alpha = dnnType(1), beta = dnnType(1); // place bias into dstData - checkCuda( cudaMemcpy(dstData, bias_d, dim_y*sizeof(value_type), cudaMemcpyDeviceToDevice) ); + checkCuda( cudaMemcpy(dstData, bias_d, dim_y*sizeof(dnnType), cudaMemcpyDeviceToDevice) ); //do matrix moltiplication checkERROR( cublasSgemv(net->cublasHandle, CUBLAS_OP_T, diff --git a/src/Flatten.cpp b/src/Flatten.cpp index 7751c8d..732d941 100644 --- a/src/Flatten.cpp +++ b/src/Flatten.cpp @@ -7,7 +7,7 @@ namespace tkDNN { Flatten::Flatten(Network *net) : Layer(net) { - checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) ); + checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) ); output_dim.n = 1; output_dim.c = input_dim.tot(); @@ -22,7 +22,7 @@ Flatten::~Flatten() { checkCuda( cudaFree(dstData) ); } -value_type* Flatten::infer(dataDim_t &dim, value_type* srcData) { +dnnType* Flatten::infer(dataDim_t &dim, dnnType* srcData) { //transpose per channel matrixTranspose(net->cublasHandle, srcData, dstData, dim.c, dim.h*dim.w*dim.l); diff --git a/src/MulAdd.cpp b/src/MulAdd.cpp index 9d9745b..2e4a824 100644 --- a/src/MulAdd.cpp +++ b/src/MulAdd.cpp @@ -5,7 +5,7 @@ namespace tkDNN { -MulAdd::MulAdd(Network *net, value_type mul, value_type add) : Layer(net) { +MulAdd::MulAdd(Network *net, dnnType mul, dnnType add) : Layer(net) { this->mul = mul; this->add = add; @@ -13,16 +13,16 @@ MulAdd::MulAdd(Network *net, value_type mul, value_type add) : Layer(net) { int size = input_dim.tot(); // create a vector with all value setted to add - value_type *add_vector_h = new value_type[size]; + dnnType *add_vector_h = new dnnType[size]; for(int i=0; icublasHandle, srcData, dstData, add_vector, input_dim.tot(), mul); diff --git a/src/Network.cpp b/src/Network.cpp index 61153df..fa97c25 100644 --- a/src/Network.cpp +++ b/src/Network.cpp @@ -30,7 +30,7 @@ Network::~Network() { checkERROR( cublasDestroy(cublasHandle) ); } -value_type* Network::infer(dataDim_t &dim, value_type* data) { +dnnType* Network::infer(dataDim_t &dim, dnnType* data) { //do infer for every layer for(int i=0; i output index = "<enqueue(1, buffersRT, stream, nullptr); @@ -161,7 +161,7 @@ ITensor* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { float eps = CUDNN_BN_MIN_EPSILON; //make power array of ones - value_type *power_h = new value_type[l->outputs]; + dnnType *power_h = new dnnType[l->outputs]; for(int i=0; ioutputs; i++) power_h[i] = 1.0f; //convert mean diff --git a/src/Pooling.cpp b/src/Pooling.cpp index e4cacbd..c66829e 100644 --- a/src/Pooling.cpp +++ b/src/Pooling.cpp @@ -57,15 +57,15 @@ Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW, output_dim.w = w; output_dim.l = l; - checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) ); + checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); //pool on 3d data need transposition at the enter and on the exit //allocate for initial and final transposition if(poolOn3d) { output_dim.n = 1; - checkCuda( cudaMalloc(&tmpInputData, input_dim.tot()*sizeof(value_type)) ); - checkCuda( cudaMalloc(&tmpOutputData, output_dim.tot()*sizeof(value_type)) ); + checkCuda( cudaMalloc(&tmpInputData, input_dim.tot()*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&tmpOutputData, output_dim.tot()*sizeof(dnnType)) ); } } @@ -81,10 +81,10 @@ Pooling::~Pooling() { checkCuda( cudaFree(dstData) ); } -value_type* Pooling::infer(dataDim_t &dim, value_type* srcData) { +dnnType* Pooling::infer(dataDim_t &dim, dnnType* srcData) { - value_type *poolSrc = srcData; - value_type *poolDst = dstData; + dnnType *poolSrc = srcData; + dnnType *poolDst = dstData; if(poolOn3d) { matrixTranspose(net->cublasHandle, srcData, tmpInputData, dim.h*dim.w*dim.c, dim.l); @@ -92,8 +92,8 @@ value_type* Pooling::infer(dataDim_t &dim, value_type* srcData) { poolDst = tmpOutputData; } - value_type alpha = value_type(1); - value_type beta = value_type(0); + dnnType alpha = dnnType(1); + dnnType beta = dnnType(0); checkCUDNN( cudnnPoolingForward(net->cudnnHandle, poolingDesc, &alpha, srcTensorDesc, poolSrc, &beta, dstTensorDesc, poolDst) ); diff --git a/src/Region.cpp b/src/Region.cpp index 2fe0039..1861488 100644 --- a/src/Region.cpp +++ b/src/Region.cpp @@ -1,5 +1,10 @@ #include +#ifdef OPENCV + #include + #include +#endif + #include "Layer.h" #include "kernels.h" @@ -12,6 +17,7 @@ Region::Region(Network *net, int classes, int coords, int num, float thresh, con this->coords = coords; this->num = num; this->thresh = thresh; + this->res_boxes_n = 0; // same output_dim.n = input_dim.n; @@ -23,7 +29,7 @@ Region::Region(Network *net, int classes, int coords, int num, float thresh, con //load anchors readBinaryFile(fname_weights, 2*num, &bias_h, &bias_d); - checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) ); + checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) ); } Region::~Region() { @@ -36,9 +42,9 @@ int Region::entry_index(int batch, int location, int entry) { return batch*output_dim.tot() + n*input_dim.w*input_dim.h*(coords+classes+1) + entry*input_dim.w*input_dim.h + loc; } -value_type* Region::infer(dataDim_t &dim, value_type* srcData) { +dnnType* Region::infer(dataDim_t &dim, dnnType* srcData) { - checkCuda( cudaMemcpy(dstData, srcData, dim.tot()*sizeof(value_type), cudaMemcpyDeviceToDevice)); + checkCuda( cudaMemcpy(dstData, srcData, dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice)); for (int b = 0; b < dim.n; ++b){ for(int n = 0; n < num; ++n){ @@ -192,11 +198,11 @@ int max_index(float *a, int n) { void Region::interpretData() { - int imW = 768, imH = 576; + int imW = net->input_dim.w, imH = net->input_dim.h; int tot = output_dim.w*output_dim.h*num; - float *lel = new value_type[output_dim.tot()]; - cudaMemcpy(lel, dstData, output_dim.tot()*sizeof(value_type), cudaMemcpyDeviceToHost); + 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 *)); @@ -232,10 +238,51 @@ void Region::interpretData() { int cl = max_index(probs[i], classes); float prob = probs[i][cl]; if(prob > thresh) { - //printf("%d %s: %.0f%%\n", i, names[class], prob*100); - printf("%d: %.0f%%\n", cl, prob*100); + box b = boxes[i]; + int x = (b.x-b.w/2.)*imW; + int w = (b.x+b.w/2.)*imW - b.x; + int y = (b.y-b.h/2.)*imH; + int h = (b.y+b.h/2.)*imH - b.y; + + printf("%d: %.0f%% box(x1, y1, x2, y2): %d %d %d %d\n", cl, prob*100, x, y, w, h); + b.x = x; + b.y = y; + b.h = h; + b.w = w; + res_boxes[res_boxes_n] = b; + res_boxes_n++; } } } +void Region::showImageResult(dnnType *input_h) { + +#ifdef OPENCV + dataDim_t dim = net->input_dim; + // read an image + cv::Mat r(dim.h, dim.w, CV_32F, input_h); + cv::Mat g(dim.h, dim.w, CV_32F, input_h + dim.h*dim.w); + cv::Mat b(dim.h, dim.w, CV_32F, input_h + dim.h*dim.w*2); + std::vector array_to_merge; + array_to_merge.push_back(b); + array_to_merge.push_back(g); + array_to_merge.push_back(r); + cv::Mat color; + cv::merge(array_to_merge, color); + + for(int i=0; idstData; + dnnType *input = layers[i]->dstData; int in_dim = layers[i]->input_dim.tot(); - checkCuda( cudaMemcpy(dstData + offset, input, in_dim*sizeof(value_type), cudaMemcpyDeviceToDevice)); + checkCuda( cudaMemcpy(dstData + offset, input, in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); offset += in_dim; } diff --git a/src/Softmax.cpp b/src/Softmax.cpp index f3b85fe..076b723 100644 --- a/src/Softmax.cpp +++ b/src/Softmax.cpp @@ -7,7 +7,7 @@ namespace tkDNN { Softmax::Softmax(Network *net) : Layer(net) { - checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) ); + checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) ); checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc, net->tensorFormat, @@ -28,10 +28,10 @@ Softmax::~Softmax() { checkCuda( cudaFree(dstData) ); } -value_type* Softmax::infer(dataDim_t &dim, value_type* srcData) { +dnnType* Softmax::infer(dataDim_t &dim, dnnType* srcData) { - value_type alpha = value_type(1); - value_type beta = value_type(0); + dnnType alpha = dnnType(1); + dnnType beta = dnnType(0); checkCUDNN( cudnnSoftmaxForward(net->cudnnHandle, CUDNN_SOFTMAX_ACCURATE , CUDNN_SOFTMAX_MODE_CHANNEL, diff --git a/src/kernels/activation_elu.cu b/src/kernels/activation_elu.cu index 073d5cf..cb6d420 100644 --- a/src/kernels/activation_elu.cu +++ b/src/kernels/activation_elu.cu @@ -7,12 +7,12 @@ x > 0 : y = x */ __global__ -void activation_elu(value_type *input, value_type *output, int size) { +void activation_elu(dnnType *input, dnnType *output, int size) { int i = blockDim.x*blockIdx.x + threadIdx.x; if(i0) k0 = 1.0f; @@ -28,7 +28,7 @@ void activation_elu(value_type *input, value_type *output, int size) { /** ELU activation function */ -void activationELUForward(value_type* srcData, value_type* dstData, int size) +void activationELUForward(dnnType* srcData, dnnType* dstData, int size) { int blocks = (size+255)/256; int threads = 256; diff --git a/src/kernels/activation_leaky.cu b/src/kernels/activation_leaky.cu index d22ca0c..60466f7 100644 --- a/src/kernels/activation_leaky.cu +++ b/src/kernels/activation_leaky.cu @@ -1,7 +1,7 @@ #include "kernels.h" __global__ -void activation_leaky(value_type *input, value_type *output, int size) { +void activation_leaky(dnnType *input, dnnType *output, int size) { int i = blockDim.x*blockIdx.x + threadIdx.x; @@ -17,7 +17,7 @@ void activation_leaky(value_type *input, value_type *output, int size) { /** ELU activation function */ -void activationLEAKYForward(value_type* srcData, value_type* dstData, int size) +void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size) { int blocks = (size+255)/256; int threads = 256; diff --git a/src/kernels/activation_logistic.cu b/src/kernels/activation_logistic.cu index 23ace0c..3f59cad 100644 --- a/src/kernels/activation_logistic.cu +++ b/src/kernels/activation_logistic.cu @@ -1,7 +1,7 @@ #include "kernels.h" __global__ -void activation_logistic(value_type *input, value_type *output, int size) { +void activation_logistic(dnnType *input, dnnType *output, int size) { int i = blockDim.x*blockIdx.x + threadIdx.x; @@ -14,7 +14,7 @@ void activation_logistic(value_type *input, value_type *output, int size) { /** LOGISTIC activation function */ -void activationLOGISTICForward(value_type* srcData, value_type* dstData, int size) +void activationLOGISTICForward(dnnType* srcData, dnnType* dstData, int size) { int blocks = (size+255)/256; int threads = 256; diff --git a/src/kernels/reorg.cu b/src/kernels/reorg.cu index 2b7cca0..a59ec2e 100644 --- a/src/kernels/reorg.cu +++ b/src/kernels/reorg.cu @@ -35,7 +35,7 @@ __global__ void reorg_kernel(int N, float *x, int w, int h, int c, int batch, in /** reorg function function */ -void reorgForward(value_type* srcData, value_type* dstData, +void reorgForward(dnnType* srcData, dnnType* dstData, int n, int c, int h, int w, int stride) { int size = n*c*h*w; diff --git a/src/pluginsRT/ActivationLeakyRT.cpp b/src/pluginsRT/ActivationLeakyRT.cpp index 355d8ea..c693730 100644 --- a/src/pluginsRT/ActivationLeakyRT.cpp +++ b/src/pluginsRT/ActivationLeakyRT.cpp @@ -41,8 +41,8 @@ public: virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { - activationLEAKYForward((value_type*)reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), size); + activationLEAKYForward((dnnType*)reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), size); return 0; } diff --git a/src/pluginsRT/RegionRT.cpp b/src/pluginsRT/RegionRT.cpp index fb5a82d..844199e 100644 --- a/src/pluginsRT/RegionRT.cpp +++ b/src/pluginsRT/RegionRT.cpp @@ -44,10 +44,10 @@ public: virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { - value_type *srcData = (value_type*)reinterpret_cast(inputs[0]); - value_type *dstData = reinterpret_cast(outputs[0]); + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); - checkCuda( cudaMemcpy(dstData, srcData, batchSize*c*h*w*sizeof(value_type), cudaMemcpyDeviceToDevice)); + checkCuda( cudaMemcpy(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice)); for (int b = 0; b < batchSize; ++b){ for(int n = 0; n < num; ++n){ diff --git a/src/pluginsRT/ReorgRT.cpp b/src/pluginsRT/ReorgRT.cpp index 9d8014d..766c33a 100644 --- a/src/pluginsRT/ReorgRT.cpp +++ b/src/pluginsRT/ReorgRT.cpp @@ -40,8 +40,8 @@ public: virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { - reorgForward((value_type*)reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), + reorgForward((dnnType*)reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize, c, h, w, stride); return 0; } diff --git a/src/utils.cpp b/src/utils.cpp index f3e765f..8e2a12d 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -14,7 +14,7 @@ void printCenteredTitle(const char *title, char fill, int dim) { } -void readBinaryFile(const char* fname, int size, value_type** data_h, value_type** data_d, int seek) +void readBinaryFile(const char* fname, int size, dnnType** data_h, dnnType** data_d, int seek) { std::ifstream dataFile (fname, std::ios::in | std::ios::binary); std::stringstream error_s; @@ -25,11 +25,11 @@ void readBinaryFile(const char* fname, int size, value_type** data_h, value_type } if(seek != 0) { - dataFile.seekg(seek*sizeof(value_type), dataFile.cur); + dataFile.seekg(seek*sizeof(dnnType), dataFile.cur); } - int size_b = size*sizeof(value_type); - *data_h = new value_type[size]; + int size_b = size*sizeof(dnnType); + *data_h = new dnnType[size]; if (!dataFile.read ((char*) *data_h, size_b)) { error_s << "Error reading file " << fname; @@ -40,13 +40,13 @@ void readBinaryFile(const char* fname, int size, value_type** data_h, value_type checkCuda( cudaMemcpy(*data_d, *data_h, size_b, cudaMemcpyHostToDevice) ); } -void printDeviceVector(int size, value_type* vec_d, bool device) +void printDeviceVector(int size, dnnType* vec_d, bool device) { - value_type *vec; + dnnType *vec; if(device) { - vec = new value_type[size]; + vec = new dnnType[size]; cudaDeviceSynchronize(); - cudaMemcpy(vec, vec_d, size*sizeof(value_type), cudaMemcpyDeviceToHost); + cudaMemcpy(vec, vec_d, size*sizeof(dnnType), cudaMemcpyDeviceToHost); } else { vec = vec_d; } @@ -60,17 +60,17 @@ void printDeviceVector(int size, value_type* vec_d, bool device) delete [] vec; } -int checkResult(int size, value_type *data_d, value_type *correct_d, bool device) { +int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) { - value_type *data_h, *correct_h; + dnnType *data_h, *correct_h; const float eps = 0.0001f; if(device) { - data_h = new value_type[size]; - correct_h = new value_type[size]; + data_h = new dnnType[size]; + correct_h = new dnnType[size]; cudaDeviceSynchronize(); - cudaMemcpy(data_h, data_d, size*sizeof(value_type), cudaMemcpyDeviceToHost); - cudaMemcpy(correct_h, correct_d, size*sizeof(value_type), cudaMemcpyDeviceToHost); + cudaMemcpy(data_h, data_d, size*sizeof(dnnType), cudaMemcpyDeviceToHost); + cudaMemcpy(correct_h, correct_d, size*sizeof(dnnType), cudaMemcpyDeviceToHost); } else { data_h = data_d; @@ -103,30 +103,30 @@ int checkResult(int size, value_type *data_d, value_type *correct_d, bool device return diffs; } -void resize(int size, value_type **data) +void resize(int size, dnnType **data) { if (*data != NULL) checkCuda( cudaFree(*data) ); - checkCuda( cudaMalloc(data, size*sizeof(value_type)) ); + checkCuda( cudaMalloc(data, size*sizeof(dnnType)) ); } -void matrixTranspose(cublasHandle_t handle, value_type* srcData, value_type* dstData, int rows, int cols) { +void matrixTranspose(cublasHandle_t handle, dnnType* srcData, dnnType* dstData, int rows, int cols) { - value_type *A = srcData, *clone = dstData; + dnnType *A = srcData, *clone = dstData; int m = rows, n= cols; - checkCuda( cudaMemcpy(clone, A, m*n*sizeof(value_type), cudaMemcpyDeviceToDevice)); + checkCuda( cudaMemcpy(clone, A, m*n*sizeof(dnnType), cudaMemcpyDeviceToDevice)); float const alpha(1.0); float const beta(0.0); checkERROR( cublasSgeam( handle, CUBLAS_OP_T, CUBLAS_OP_N, m, n, &alpha, A, n, &beta, A, m, clone, m )); } -void matrixMulAdd( cublasHandle_t handle, value_type* srcData, value_type* dstData, - value_type* add_vector, int dim, value_type mul) { +void matrixMulAdd( cublasHandle_t handle, dnnType* srcData, dnnType* dstData, + dnnType* add_vector, int dim, dnnType mul) { - checkCuda( cudaMemcpy(dstData, add_vector, dim*sizeof(value_type), cudaMemcpyDeviceToDevice)); + checkCuda( cudaMemcpy(dstData, add_vector, dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); - value_type alpha = mul; + dnnType alpha = mul; checkERROR( cublasSaxpy(handle, dim, &alpha, srcData, 1, dstData, 1)); } diff --git a/tests/mnist/test_mnist.cpp b/tests/mnist/test_mnist.cpp index 44b2dbb..747b2d7 100644 --- a/tests/mnist/test_mnist.cpp +++ b/tests/mnist/test_mnist.cpp @@ -25,11 +25,11 @@ int main() { tkDNN::NetworkRT netRT(&net); // Load input - value_type *data; - value_type *input_h; + dnnType *data; + dnnType *input_h; readBinaryFile(input_bin, dim.tot(), &input_h, &data); - value_type *out_data, *out_data2; + dnnType *out_data, *out_data2; std::cout<<"CUDNN inference:\n"; { dim.print(); //print initial dimension @@ -63,8 +63,8 @@ int main() { /* // Print real test std::cout<<"\n==== CHECK RESULT ====\n"; - value_type *out; - value_type *out_h; + dnnType *out; + dnnType *out_h; readBinaryFile(output_bin, dim.tot(), &out_h, &out); printDeviceVector(dim.tot(), out); */ diff --git a/tests/mnist/test_mnistRT.cpp b/tests/mnist/test_mnistRT.cpp index 74a2d2f..a0814c6 100644 --- a/tests/mnist/test_mnistRT.cpp +++ b/tests/mnist/test_mnistRT.cpp @@ -39,8 +39,8 @@ int main() { tkDNN::Softmax l7(&net); // Load input - value_type *data; - value_type *input_h; + dnnType *data; + dnnType *input_h; readBinaryFile(input_bin, dim.tot(), &input_h, &data); dim.print(); //print initial dimension @@ -55,8 +55,8 @@ int main() { // Print real test std::cout<<"\n==== CHECK CUDNN RESULT ====\n"; - value_type *out; - value_type *out_h; + dnnType *out; + dnnType *out_h; readBinaryFile(output_bin, dim.tot(), &out_h, &out); std::cout<<"Diff: "<