Improve Yolo, mobilenet and shelfnet preprocessing using GPU #56

Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>

include/tkDNN/DetectionNN.h

@@ -1,6 +1,8 @@
 #ifndef DETECTIONNN_H
 #define DETECTIONNN_H
 
+#include "kernels.h"
+
 #include <iostream>
 #include <signal.h>
 #include <stdlib.h>    
@@ -16,6 +18,7 @@
 #include <opencv2/imgproc/imgproc.hpp>
 
 #include "tkdnn.h"
+#include "utilsNN.h"
 
 //#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
 
@@ -31,6 +34,7 @@ class DetectionNN {
 
     protected:
         tk::dnn::NetworkRT *netRT = nullptr;
+        uint8_t *frame_d = nullptr;
         dnnType *input_d;
 
         std::vector<cv::Size> originalSize;
@@ -38,6 +42,7 @@ class DetectionNN {
         cv::Scalar colors[256];
 
         int nBatches = 1;
+        int frame_size = 0;
 
 #ifdef OPENCV_CUDACONTRIB
         cv::cuda::GpuMat bgr[3];

include/tkDNN/SegmentationNN.h

@@ -1,6 +1,8 @@
 #ifndef SEGMENTATIONNN_H
 #define SEGMENTATIONNN_H
 
+#include "kernels.h"
+
 #include <iostream>
 #include <signal.h>
 #include <stdlib.h>    
@@ -16,6 +18,7 @@
 #include "tkdnn.h"
 #include "NetworkViz.h"
 #include "kernelsThrust.h"
+#include "utilsNN.h"
 
 namespace tk { namespace dnn {
 
@@ -23,7 +26,9 @@ class SegmentationNN {
 
     protected:
         tk::dnn::NetworkRT *netRT = nullptr;
+        uint8_t *frame_d = nullptr;
         int nBatches = 1;
+        int frame_size = 0;
 
         std::vector<cv::Size> originalSize;
         cv::Mat bgr[3];
@@ -76,18 +81,7 @@ class SegmentationNN {
             cv::copyMakeBorder(frame, frame_cropped, top, bottom, left, right, cv::BORDER_CONSTANT, cv::Scalar(0,0,0) );
 
             tk::dnn::dataDim_t idim = netRT->input_dim;
-
-            resize(frame_cropped, frame_cropped, cv::Size(idim.w, idim.h));
-            
-            cv::split(frame_cropped, bgr);
-            for (int i = 0; i < idim.c; i++){
-                int idx = i * frame_cropped.rows * frame_cropped.cols;
-                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+ idim.tot()*bi, input + idim.tot()*bi, idim.tot() * sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
-
+            resizeAndSplit(frame_cropped, &frame_d, frame_size, input_d, netRT, bi, true);
             normalize(input_d + idim.tot()*bi, idim.c, idim.h, idim.w, mean_d, stddev_d);
         }        
 

include/tkDNN/kernels.h

@@ -48,4 +48,11 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
                          const int dst_dim, cudaStream_t stream = cudaStream_t(0));
 
 void scalAdd(dnnType* dstData, int size, float alpha, float beta, int inc, cudaStream_t stream = cudaStream_t(0));
+
+void normalize(float *bgr, const int ch, const int h, const int w, const float *mean, const float *stddev);
+void normalize(float *bgr, const int ch, const int h, const int w, const float mean, const float stddev);
+
+void interleavedToPlanar( uint8_t *d_src, float *d_dst, int s_w, int s_h, int s_c, int d_w, int d_h);
+void interleavedRGBToPlanarBGR( uint8_t *d_src, float *d_dst, int s_w, int s_h, int s_c, int d_w, int d_h);
+
 #endif //KERNELS_H

include/tkDNN/kernelsThrust.h

@@ -31,7 +31,6 @@ 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 normalize(float *bgr, const int ch, const int h, const int w, const float *mean, const float *stddev);
 void transformDep(float *src_begin, float *src_end, float *dst_begin, float *dst_end);
 void subtractWithThreshold(dnnType *src_begin, dnnType *src_end, dnnType *src2_begin, dnnType *src_out, struct threshold op);
 void topKxyclasses(int *ids_begin, int *ids_end, const int K, const int size, const int wh, int *clses, int *xs, int *ys);

include/tkDNN/utilsNN.h

@@ -0,0 +1,10 @@
+#ifndef UTILSNN_H
+#define UTILSNN_H
+
+#include "tkdnn.h"
+#include <opencv2/core/core.hpp>
+
+
+void resizeAndSplit(cv::Mat& frame, uint8_t** frame_d, int& frame_size, dnnType *input_d, tk::dnn::NetworkRT *netRT, const int bi=0, bool BGR=true);
+
+#endif // UTILSNN_H

src/MobilenetDetection.cpp

@@ -211,38 +211,8 @@ bool MobilenetDetection::init(const std::string& tensor_path, const int n_classe
 }
 
 void MobilenetDetection::preprocess(cv::Mat &frame, const int bi){
-#ifdef OPENCV_CUDACONTRIB
-        //move original image on GPU
-        cv::cuda::GpuMat orig_img, frame_nomean;
-        orig_img = cv::cuda::GpuMat(frame);
-
-        //resize image, remove mean, divide by std
-        cv::cuda::resize (orig_img, orig_img, cv::Size(netRT->input_dim.w, netRT->input_dim.h)); 
-        orig_img.convertTo(frame_nomean, CV_32FC3, 1, -127);
-        frame_nomean.convertTo(imagePreproc, CV_32FC3, 1 / 128.0, 0);
-
-        //copy image into tensors
-        cv::cuda::split(imagePreproc, bgr);
-
-        for(int i=0; i < netRT->input_dim.c; i++){
-            int idx = i * imagePreproc.rows * imagePreproc.cols;
-            checkCuda( cudaMemcpy((void *)&input_d[idx + netRT->input_dim.tot()*bi], (void *)bgr[i].data, imagePreproc.rows * imagePreproc.cols* sizeof(float), cudaMemcpyDeviceToDevice) );
-        }
-#else
-        //resize image, remove mean, divide by std
-        cv::Mat frame_nomean;
-        resize(frame, frame, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
-        frame.convertTo(frame_nomean, CV_32FC3, 1, -127);
-        frame_nomean.convertTo(imagePreproc, CV_32FC3, 1 / 128.0, 0);
-
-        //copy image into tensor and copy it into GPU
-        cv::split(imagePreproc, bgr);
-        for (int i = 0; i < netRT->input_dim.c; i++){
-            int idx = i * imagePreproc.rows * imagePreproc.cols;
-            memcpy((void *)&input[idx + netRT->input_dim.tot()*bi], (void *)bgr[i].data, imagePreproc.rows * imagePreproc.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));
-#endif
+    resizeAndSplit(frame, &frame_d, frame_size, input_d, netRT, bi, false);
+    normalize(input_d +  netRT->input_dim.tot()*bi,  netRT->input_dim.c,  netRT->input_dim.h,  netRT->input_dim.w, 127.0f, 128.0f);
 }
 
 void MobilenetDetection::postprocess(const int bi, const bool mAP){

src/Yolo3Detection.cpp

@@ -57,38 +57,8 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, c
 } 
 
 void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){
-#ifdef OPENCV_CUDACONTRIB
-    cv::cuda::GpuMat orig_img, img_resized;
-    orig_img = cv::cuda::GpuMat(frame);
-    cv::cuda::resize(orig_img, img_resized, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
-
-    img_resized.convertTo(imagePreproc, CV_32FC3, 1/255.0); 
-
-    //split channels
-    cv::cuda::split(imagePreproc,bgr);//split source
-
-    //write channels
-    for(int i=0; i<netRT->input_dim.c; i++) {
-        int size = imagePreproc.rows * imagePreproc.cols;
-        int ch = netRT->input_dim.c-1 -i;
-        bgr[ch].download(bgr_h); //TODO: don't copy back on CPU
-        checkCuda( cudaMemcpy(input_d + i*size + netRT->input_dim.tot()*bi, (float*)bgr_h.data, size*sizeof(dnnType), cudaMemcpyHostToDevice));
-    }
-#else
-    cv::resize(frame, frame, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
-    frame.convertTo(imagePreproc, CV_32FC3, 1/255.0); 
-
-    //split channels
-    cv::split(imagePreproc,bgr);//split source
-
-    //write channels
-    for(int i=0; i<netRT->input_dim.c; i++) {
-        int idx = i*imagePreproc.rows*imagePreproc.cols;
-        int ch = netRT->input_dim.c-1 -i;
-        memcpy((void*)&input[idx + netRT->input_dim.tot()*bi], (void*)bgr[ch].data, imagePreproc.rows*imagePreproc.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));
-#endif
+    resizeAndSplit(frame, &frame_d, frame_size, input_d, netRT, bi, true);
+    normalize(input_d +  netRT->input_dim.tot()*bi,  netRT->input_dim.c,  netRT->input_dim.h,  netRT->input_dim.w, 0.0f, 255.0f);
 }
 
 void Yolo3Detection::postprocess(const int bi, const bool mAP){

src/kernels/interleaved_to_planar.cu

@@ -0,0 +1,84 @@
+#include "kernels.h"
+
+
+__global__ void interleavedToPlanarKernel(uint8_t *src, float *dst, int s_w, int s_h, int s_c, int d_w, int d_h, float ratio_w, float ratio_h) {
+
+    int x = min( (int)(blockIdx.x * blockDim.x + threadIdx.x), d_w-1);
+    int y = min( (int)(blockIdx.y * blockDim.y + threadIdx.y), d_h-1);
+
+    float sum_r=0, sum_g=0, sum_b=0;
+
+    float x_src = (float) x * ratio_w; // + ratio_w/2;
+    float y_src = (float) y * ratio_h; // + ratio_h/2;
+
+    int r = (int) y_src;
+    int c = (int) x_src;
+    float dr = y_src - r;
+    float dc = x_src - c;
+
+    sum_r = (float) src[(r * s_w + c) * s_c] * (1.0f - dr) * (1.0f - dc) +
+            (float) src[((r + 1) * s_w + c) * s_c] * (dr) * (1.0f - dc) +
+            (float) src[(r * s_w + c + 1) * s_c] * (1.0f - dr) * (dc) +
+            (float) src[((r + 1) * s_w + c + 1) * s_c] * (dr) * (dc);
+    sum_g = (float) src[(r * s_w + c) * s_c + 1] * (1.0f - dr) * (1.0f - dc) +
+            (float) src[((r + 1) * s_w + c) * s_c + 1] * (dr) * (1.0f - dc) +
+            (float) src[(r * s_w + c + 1) * s_c + 1] * (1.0f - dr) * (dc) +
+            (float) src[((r + 1) * s_w + c + 1) * s_c + 1] * (dr) * (dc);
+    sum_b = (float) src[(r * s_w + c) * s_c + 2] * (1.0f - dr) * (1.0f - dc) +
+            (float) src[((r + 1) * s_w + c) * s_c + 2] * (dr) * (1.0f - dc) +
+            (float) src[(r * s_w + c + 1) * s_c + 2] * (1.0f - dr) * (dc) +
+            (float) src[((r + 1) * s_w + c + 1) * s_c + 2] * (dr) * (dc);
+
+    dst[y * d_w + x] 					= sum_r;
+    dst[y * d_w + x + d_w * d_h] 		= sum_g;
+    dst[y * d_w + x + d_w * d_h * 2] 	= sum_b;
+}
+
+__global__ void interleavedRGBToPlanarBGRKernel(uint8_t *src, float *dst, int s_w, int s_h, int s_c, int d_w, int d_h, float ratio_w, float ratio_h) {
+
+    int x = min( (int)(blockIdx.x * blockDim.x + threadIdx.x), d_w-1);
+    int y = min( (int)(blockIdx.y * blockDim.y + threadIdx.y), d_h-1);
+
+    float sum_r=0, sum_g=0, sum_b=0;
+
+    float x_src = (float) x * ratio_w; // + ratio_w/2;
+    float y_src = (float) y * ratio_h; // + ratio_h/2;
+
+    int r = (int) y_src;
+    int c = (int) x_src;
+    float dr = y_src - r;
+    float dc = x_src - c;
+
+    sum_r = (float) src[(r * s_w + c) * s_c] * (1.0f - dr) * (1.0f - dc) +
+            (float) src[((r + 1) * s_w + c) * s_c] * (dr) * (1.0f - dc) +
+            (float) src[(r * s_w + c + 1) * s_c] * (1.0f - dr) * (dc) +
+            (float) src[((r + 1) * s_w + c + 1) * s_c] * (dr) * (dc);
+    sum_g = (float) src[(r * s_w + c) * s_c + 1] * (1.0f - dr) * (1.0f - dc) +
+            (float) src[((r + 1) * s_w + c) * s_c + 1] * (dr) * (1.0f - dc) +
+            (float) src[(r * s_w + c + 1) * s_c + 1] * (1.0f - dr) * (dc) +
+            (float) src[((r + 1) * s_w + c + 1) * s_c + 1] * (dr) * (dc);
+    sum_b = (float) src[(r * s_w + c) * s_c + 2] * (1.0f - dr) * (1.0f - dc) +
+            (float) src[((r + 1) * s_w + c) * s_c + 2] * (dr) * (1.0f - dc) +
+            (float) src[(r * s_w + c + 1) * s_c + 2] * (1.0f - dr) * (dc) +
+            (float) src[((r + 1) * s_w + c + 1) * s_c + 2] * (dr) * (dc);
+
+    dst[y * d_w + x] 					= sum_b;
+    dst[y * d_w + x + d_w * d_h] 		= sum_g;
+    dst[y * d_w + x + d_w * d_h * 2] 	= sum_r;
+}
+
+void interleavedToPlanar( uint8_t *d_src, float *d_dst, int s_w, int s_h, int s_c, int d_w, int d_h){
+	dim3 dg( ceil( (double)d_w/32 ), ceil( (double)d_h/8 ) );
+	dim3 db( 32, 8);
+
+	interleavedToPlanarKernel<<< dg, db >>>(d_src, d_dst, s_w, s_h, s_c, d_w, d_h, (float)s_w/d_w, (float)s_h/d_h);
+	cudaDeviceSynchronize();
+}
+
+void interleavedRGBToPlanarBGR( uint8_t *d_src, float *d_dst, int s_w, int s_h, int s_c, int d_w, int d_h){
+	dim3 dg( ceil( (double)d_w/32 ), ceil( (double)d_h/8 ) );
+	dim3 db( 32, 8);
+
+	interleavedRGBToPlanarBGRKernel<<< dg, db >>>(d_src, d_dst, s_w, s_h, s_c, d_w, d_h, (float)s_w/d_w, (float)s_h/d_h);
+	cudaDeviceSynchronize();
+}

src/kernels/normalize.cu

@@ -1,4 +1,4 @@
-#include "kernelsThrust.h"
+#include "kernels.h"
 
 __global__
 void normalize_kernel(float *bgr, const int dim, const float *mean, const float *stddev){
@@ -9,8 +9,25 @@ void normalize_kernel(float *bgr, const int dim, const float *mean, const float
     
 }
 
+
+__global__
+void normalize_kernel(float *bgr, const int dim, const float mean, const float stddev){
+    int i = blockDim.x*blockIdx.x + threadIdx.x;
+    int j = blockIdx.y;
+    bgr[j*(dim)+i] = bgr[j*(dim)+i] - mean;
+    bgr[j*(dim)+i] = bgr[j*(dim)+i] / stddev;
+    
+}
+
 void normalize(float *bgr, const int ch, const int h, const int w, const float *mean, const float *stddev){
     int num_thread = 256;
     dim3 dimBlock(h*w/num_thread, ch);
     normalize_kernel<<<dimBlock, num_thread, 0>>>(bgr, h*w, mean, stddev);
 }
+
+
+void normalize(float *bgr, const int ch, const int h, const int w, const float mean, const float stddev){
+    int num_thread = 256;
+    dim3 dimBlock(h*w/num_thread, ch);
+    normalize_kernel<<<dimBlock, num_thread, 0>>>(bgr, h*w, mean, stddev);
+}

src/utilsNN.cpp

@@ -0,0 +1,41 @@
+#include "kernels.h"
+#include "utilsNN.h"
+
+
+void resizeAndSplit(cv::Mat& frame, uint8_t** frame_d, int& frame_size, dnnType *input_d, tk::dnn::NetworkRT *netRT, const int bi, bool BGR){
+    int new_frame_size = sizeof(uint8_t) * frame.cols * frame.rows * frame.channels();
+    if(*frame_d == nullptr){
+        frame_size = new_frame_size;
+        checkCuda(cudaMalloc(frame_d, frame_size));
+    }
+    else{
+        if(new_frame_size > frame_size ){
+            frame_size = new_frame_size;
+            checkCuda(cudaFree(frame_d));
+            checkCuda(cudaMalloc(frame_d, frame_size));
+        }
+    }
+    
+    checkCuda(cudaMemcpyAsync(*frame_d, frame.data, frame_size, cudaMemcpyHostToDevice, netRT->stream));
+    if(BGR){
+        interleavedRGBToPlanarBGR(*frame_d,
+                                input_d + netRT->input_dim.tot() * bi, 
+                                frame.cols, 
+                                frame.rows, 
+                                frame.channels(),
+                                netRT->input_dim.w,
+                                netRT->input_dim.h
+                            );
+    }
+    else{
+        interleavedToPlanar(*frame_d,
+                                input_d + netRT->input_dim.tot() * bi, 
+                                frame.cols, 
+                                frame.rows, 
+                                frame.channels(),
+                                netRT->input_dim.w,
+                                netRT->input_dim.h
+                            );
+
+    }
+}