mmr/tkDNN
Archived
1
0
Fork 0
Code Issues 46 Pull Requests 12 Releases 4 Activity

Add interleavedToPlanar kernels, normalize kernels and utiliy funcs

Browse Source 
Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>
This commit is contained in:
Micaela Verucchi
2022-01-27 10:47:13 +01:00
parent 04de9908a6
commit c2825cc570
6 changed files with 160 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/tkDNN/kernels.h
+7
View File
@@ -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
-1
View File
@@ -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
+10
View File
@@ -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/kernels/interleaved_to_planar.cu
+84
View File
@@ -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
+18 -1
View File
@@ -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
+41
View File
@@ -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
);
}
}