From c2825cc570ac61b1c5dcc0b093938a6402e490c1 Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Thu, 27 Jan 2022 10:47:13 +0100 Subject: [PATCH] Add interleavedToPlanar kernels, normalize kernels and utiliy funcs Signed-off-by: Micaela Verucchi --- include/tkDNN/kernels.h | 7 +++ include/tkDNN/kernelsThrust.h | 1 - include/tkDNN/utilsNN.h | 10 ++++ src/kernels/interleaved_to_planar.cu | 84 ++++++++++++++++++++++++++++ src/kernels/normalize.cu | 19 ++++++- src/utilsNN.cpp | 41 ++++++++++++++ 6 files changed, 160 insertions(+), 2 deletions(-) create mode 100644 include/tkDNN/utilsNN.h create mode 100644 src/kernels/interleaved_to_planar.cu create mode 100644 src/utilsNN.cpp diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index d809129..95080a7 100644 --- a/include/tkDNN/kernels.h +++ b/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 diff --git a/include/tkDNN/kernelsThrust.h b/include/tkDNN/kernelsThrust.h index ab2fd90..fb7696a 100644 --- a/include/tkDNN/kernelsThrust.h +++ b/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); diff --git a/include/tkDNN/utilsNN.h b/include/tkDNN/utilsNN.h new file mode 100644 index 0000000..d544dcd --- /dev/null +++ b/include/tkDNN/utilsNN.h @@ -0,0 +1,10 @@ +#ifndef UTILSNN_H +#define UTILSNN_H + +#include "tkdnn.h" +#include + + +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 \ No newline at end of file diff --git a/src/kernels/interleaved_to_planar.cu b/src/kernels/interleaved_to_planar.cu new file mode 100644 index 0000000..c1a1802 --- /dev/null +++ b/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(); +} diff --git a/src/kernels/normalize.cu b/src/kernels/normalize.cu index 5206256..4511a9e 100644 --- a/src/kernels/normalize.cu +++ b/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<<>>(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<<>>(bgr, h*w, mean, stddev); } \ No newline at end of file diff --git a/src/utilsNN.cpp b/src/utilsNN.cpp new file mode 100644 index 0000000..a8dc0df --- /dev/null +++ b/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 + ); + + } +} \ No newline at end of file