Compare commits
2 Commits
| Author | SHA1 | Date | |
|---|---|---|---|
| 7c0620e391 | |||
| 3bcc32ffdc |
@@ -1,8 +1,6 @@
|
|||||||
#ifndef DETECTIONNN_H
|
#ifndef DETECTIONNN_H
|
||||||
#define DETECTIONNN_H
|
#define DETECTIONNN_H
|
||||||
|
|
||||||
#include "kernels.h"
|
|
||||||
|
|
||||||
#include <iostream>
|
#include <iostream>
|
||||||
#include <signal.h>
|
#include <signal.h>
|
||||||
#include <stdlib.h>
|
#include <stdlib.h>
|
||||||
@@ -18,7 +16,6 @@
|
|||||||
#include <opencv2/imgproc/imgproc.hpp>
|
#include <opencv2/imgproc/imgproc.hpp>
|
||||||
|
|
||||||
#include "tkdnn.h"
|
#include "tkdnn.h"
|
||||||
#include "utilsNN.h"
|
|
||||||
|
|
||||||
//#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
|
//#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
|
||||||
|
|
||||||
@@ -34,7 +31,6 @@ class DetectionNN {
|
|||||||
|
|
||||||
protected:
|
protected:
|
||||||
tk::dnn::NetworkRT *netRT = nullptr;
|
tk::dnn::NetworkRT *netRT = nullptr;
|
||||||
uint8_t *frame_d = nullptr;
|
|
||||||
dnnType *input_d;
|
dnnType *input_d;
|
||||||
|
|
||||||
std::vector<cv::Size> originalSize;
|
std::vector<cv::Size> originalSize;
|
||||||
@@ -42,7 +38,6 @@ class DetectionNN {
|
|||||||
cv::Scalar colors[256];
|
cv::Scalar colors[256];
|
||||||
|
|
||||||
int nBatches = 1;
|
int nBatches = 1;
|
||||||
int frame_size = 0;
|
|
||||||
|
|
||||||
#ifdef OPENCV_CUDACONTRIB
|
#ifdef OPENCV_CUDACONTRIB
|
||||||
cv::cuda::GpuMat bgr[3];
|
cv::cuda::GpuMat bgr[3];
|
||||||
|
|||||||
@@ -1,8 +1,6 @@
|
|||||||
#ifndef SEGMENTATIONNN_H
|
#ifndef SEGMENTATIONNN_H
|
||||||
#define SEGMENTATIONNN_H
|
#define SEGMENTATIONNN_H
|
||||||
|
|
||||||
#include "kernels.h"
|
|
||||||
|
|
||||||
#include <iostream>
|
#include <iostream>
|
||||||
#include <signal.h>
|
#include <signal.h>
|
||||||
#include <stdlib.h>
|
#include <stdlib.h>
|
||||||
@@ -18,7 +16,6 @@
|
|||||||
#include "tkdnn.h"
|
#include "tkdnn.h"
|
||||||
#include "NetworkViz.h"
|
#include "NetworkViz.h"
|
||||||
#include "kernelsThrust.h"
|
#include "kernelsThrust.h"
|
||||||
#include "utilsNN.h"
|
|
||||||
|
|
||||||
namespace tk { namespace dnn {
|
namespace tk { namespace dnn {
|
||||||
|
|
||||||
@@ -26,9 +23,7 @@ class SegmentationNN {
|
|||||||
|
|
||||||
protected:
|
protected:
|
||||||
tk::dnn::NetworkRT *netRT = nullptr;
|
tk::dnn::NetworkRT *netRT = nullptr;
|
||||||
uint8_t *frame_d = nullptr;
|
|
||||||
int nBatches = 1;
|
int nBatches = 1;
|
||||||
int frame_size = 0;
|
|
||||||
|
|
||||||
std::vector<cv::Size> originalSize;
|
std::vector<cv::Size> originalSize;
|
||||||
cv::Mat bgr[3];
|
cv::Mat bgr[3];
|
||||||
@@ -81,7 +76,18 @@ class SegmentationNN {
|
|||||||
cv::copyMakeBorder(frame, frame_cropped, top, bottom, left, right, cv::BORDER_CONSTANT, cv::Scalar(0,0,0) );
|
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;
|
tk::dnn::dataDim_t idim = netRT->input_dim;
|
||||||
resizeAndSplit(frame_cropped, &frame_d, frame_size, input_d, netRT, bi, true);
|
|
||||||
|
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));
|
||||||
|
|
||||||
normalize(input_d + idim.tot()*bi, idim.c, idim.h, idim.w, mean_d, stddev_d);
|
normalize(input_d + idim.tot()*bi, idim.c, idim.h, idim.w, mean_d, stddev_d);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|||||||
@@ -48,11 +48,4 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
|
|||||||
const int dst_dim, cudaStream_t stream = cudaStream_t(0));
|
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 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
|
#endif //KERNELS_H
|
||||||
|
|||||||
@@ -31,6 +31,7 @@ void sort(dnnType *src_begin, dnnType *src_end, int *idsrc);
|
|||||||
void topk(dnnType *src_begin, int *idsrc, int K, float *topk_scores,
|
void topk(dnnType *src_begin, int *idsrc, int K, float *topk_scores,
|
||||||
int *topk_inds, float *topk_ys, float *topk_xs);
|
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 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 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 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);
|
void topKxyclasses(int *ids_begin, int *ids_end, const int K, const int size, const int wh, int *clses, int *xs, int *ys);
|
||||||
|
|||||||
@@ -1,10 +0,0 @@
|
|||||||
#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
|
|
||||||
+43
-39
@@ -1,6 +1,6 @@
|
|||||||
#!/bin/bash
|
#!/bin/bash
|
||||||
|
|
||||||
cd build
|
#cd build
|
||||||
|
|
||||||
RED='\033[1;31m'
|
RED='\033[1;31m'
|
||||||
GREEN='\033[1;32m'
|
GREEN='\033[1;32m'
|
||||||
@@ -29,24 +29,28 @@ function print_output {
|
|||||||
|
|
||||||
}
|
}
|
||||||
|
|
||||||
|
out_dir=results
|
||||||
out_file=results.log
|
out_file=results.log
|
||||||
rm $out_file
|
rm -rf $out_dir/
|
||||||
|
mkdir -p $out_dir
|
||||||
|
|
||||||
function test_net {
|
function test_net {
|
||||||
./test_$1 &>> $out_file
|
./test_$1 &> $out_dir/$1_${TKDNN_MODE}_build_$out_file
|
||||||
print_output $? $1
|
print_output $? $1
|
||||||
./test_rtinference $1*.rt $TKDNN_BATCHSIZE &>> $out_file
|
./test_rtinference $1*.rt 1 &> $out_dir/$1_${TKDNN_MODE}_inference_batch1_$out_file
|
||||||
|
print_output $? "infer $1"
|
||||||
|
./test_rtinference $1*.rt $TKDNN_BATCHSIZE &> $out_dir/$1_${TKDNN_MODE}_inference_batch${TKDNN_BATCHSIZE}_$out_file
|
||||||
print_output $? "batched $1"
|
print_output $? "batched $1"
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
modes=( 1 ) # only FP32
|
# modes=( 1 ) # only FP32
|
||||||
# modes=( 1 2 ) # FP32 and FP16
|
modes=( 1 2 ) # FP32 and FP16
|
||||||
# modes=( 1 2 3 ) # FP32, FP16 and INT8
|
# modes=( 1 2 3 ) # FP32, FP16 and INT8
|
||||||
|
|
||||||
for i in "${modes[@]}"
|
for i in "${modes[@]}"
|
||||||
do
|
do
|
||||||
rm *rt
|
rm -f *rt
|
||||||
if [ $i -eq 1 ]
|
if [ $i -eq 1 ]
|
||||||
then
|
then
|
||||||
export TKDNN_MODE=FP32
|
export TKDNN_MODE=FP32
|
||||||
@@ -73,37 +77,37 @@ do
|
|||||||
# print_output $? imuodom
|
# print_output $? imuodom
|
||||||
|
|
||||||
test_net yolo4
|
test_net yolo4
|
||||||
test_net yolo4_320
|
# test_net yolo4_320
|
||||||
test_net yolo4_320_coco2
|
# test_net yolo4_320_coco2
|
||||||
test_net yolo4_512
|
# test_net yolo4_512
|
||||||
test_net yolo4_608
|
# test_net yolo4_608
|
||||||
test_net yolo4-csp
|
# test_net yolo4-csp
|
||||||
test_net yolo4x
|
# test_net yolo4x
|
||||||
test_net yolo4_berkeley
|
# test_net yolo4_berkeley
|
||||||
test_net yolo4_berkeley_f1
|
# test_net yolo4_berkeley_f1
|
||||||
test_net yolo4tiny
|
# test_net yolo4tiny
|
||||||
test_net yolo4tiny_512
|
# test_net yolo4tiny_512
|
||||||
test_net yolo3
|
# test_net yolo3
|
||||||
test_net yolo3_berkeley
|
# test_net yolo3_berkeley
|
||||||
test_net yolo3_coco4
|
# test_net yolo3_coco4
|
||||||
test_net yolo3_flir
|
# test_net yolo3_flir
|
||||||
test_net yolo3_512
|
# test_net yolo3_512
|
||||||
test_net yolo3tiny
|
# test_net yolo3tiny
|
||||||
test_net yolo3tiny_512
|
# test_net yolo3tiny_512
|
||||||
test_net yolo2
|
# test_net yolo2
|
||||||
test_net yolo2_voc
|
# test_net yolo2_voc
|
||||||
#test_net yolo2tiny
|
# test_net yolo2tiny
|
||||||
test_net csresnext50-panet-spp
|
# test_net csresnext50-panet-spp
|
||||||
#test_net csresnext50-panet-spp_berkeley
|
# test_net csresnext50-panet-spp_berkeley
|
||||||
test_net resnet101_cnet
|
# test_net resnet101_cnet
|
||||||
test_net dla34_cnet
|
# test_net dla34_cnet
|
||||||
test_net dla34_cnet3d
|
# test_net dla34_cnet3d
|
||||||
test_net mobilenetv2ssd
|
# test_net mobilenetv2ssd
|
||||||
test_net mobilenetv2ssd512
|
# test_net mobilenetv2ssd512
|
||||||
test_net bdd-mobilenetv2ssd
|
# test_net bdd-mobilenetv2ssd
|
||||||
test_net dla34_ctrack
|
# test_net dla34_ctrack
|
||||||
test_net shelfnet
|
# test_net shelfnet
|
||||||
test_net shelfnet_berkeley
|
# test_net shelfnet_berkeley
|
||||||
done
|
done
|
||||||
|
|
||||||
echo "If errors occured, check logfile $out_file"
|
echo "If errors occured, check logfiles in directory: $out_dir"
|
||||||
|
|||||||
@@ -211,8 +211,38 @@ bool MobilenetDetection::init(const std::string& tensor_path, const int n_classe
|
|||||||
}
|
}
|
||||||
|
|
||||||
void MobilenetDetection::preprocess(cv::Mat &frame, const int bi){
|
void MobilenetDetection::preprocess(cv::Mat &frame, const int bi){
|
||||||
resizeAndSplit(frame, &frame_d, frame_size, input_d, netRT, bi, false);
|
#ifdef OPENCV_CUDACONTRIB
|
||||||
normalize(input_d + netRT->input_dim.tot()*bi, netRT->input_dim.c, netRT->input_dim.h, netRT->input_dim.w, 127.0f, 128.0f);
|
//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
|
||||||
}
|
}
|
||||||
|
|
||||||
void MobilenetDetection::postprocess(const int bi, const bool mAP){
|
void MobilenetDetection::postprocess(const int bi, const bool mAP){
|
||||||
|
|||||||
+3
-2
@@ -278,6 +278,7 @@ void Yolo::mergeDetections(Yolo::detection *dets, int ndets, int classes, double
|
|||||||
}
|
}
|
||||||
total = k+1;
|
total = k+1;
|
||||||
|
|
||||||
|
float thresh = 0.45f;
|
||||||
for(k = 0; k < classes; ++k){
|
for(k = 0; k < classes; ++k){
|
||||||
for(i = 0; i < total; ++i){
|
for(i = 0; i < total; ++i){
|
||||||
dets[i].sort_class = k;
|
dets[i].sort_class = k;
|
||||||
@@ -288,9 +289,9 @@ void Yolo::mergeDetections(Yolo::detection *dets, int ndets, int classes, double
|
|||||||
box a = dets[i].bbox;
|
box a = dets[i].bbox;
|
||||||
for(j = i+1; j < total; ++j){
|
for(j = i+1; j < total; ++j){
|
||||||
box b = dets[j].bbox;
|
box b = dets[j].bbox;
|
||||||
if (nsm_kind == GREEDY_NMS && yolo_box_iou(a, b) > nms_thresh)
|
if (nsm_kind == GREEDY_NMS && yolo_box_iou(a, b) > thresh)
|
||||||
dets[j].prob[k] = 0;
|
dets[j].prob[k] = 0;
|
||||||
else if (nsm_kind == DIOU_NMS && yolo_box_diou(a, b, nms_thresh) > nms_thresh)
|
else if (nsm_kind == DIOU_NMS && yolo_box_diou(a, b, nms_thresh) > thresh)
|
||||||
dets[j].prob[k] = 0;
|
dets[j].prob[k] = 0;
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|||||||
+32
-2
@@ -57,8 +57,38 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, c
|
|||||||
}
|
}
|
||||||
|
|
||||||
void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){
|
void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){
|
||||||
resizeAndSplit(frame, &frame_d, frame_size, input_d, netRT, bi, true);
|
#ifdef OPENCV_CUDACONTRIB
|
||||||
normalize(input_d + netRT->input_dim.tot()*bi, netRT->input_dim.c, netRT->input_dim.h, netRT->input_dim.w, 0.0f, 255.0f);
|
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
|
||||||
}
|
}
|
||||||
|
|
||||||
void Yolo3Detection::postprocess(const int bi, const bool mAP){
|
void Yolo3Detection::postprocess(const int bi, const bool mAP){
|
||||||
|
|||||||
@@ -1,84 +0,0 @@
|
|||||||
#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();
|
|
||||||
}
|
|
||||||
@@ -1,4 +1,4 @@
|
|||||||
#include "kernels.h"
|
#include "kernelsThrust.h"
|
||||||
|
|
||||||
__global__
|
__global__
|
||||||
void normalize_kernel(float *bgr, const int dim, const float *mean, const float *stddev){
|
void normalize_kernel(float *bgr, const int dim, const float *mean, const float *stddev){
|
||||||
@@ -9,25 +9,8 @@ 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){
|
void normalize(float *bgr, const int ch, const int h, const int w, const float *mean, const float *stddev){
|
||||||
int num_thread = 256;
|
int num_thread = 256;
|
||||||
dim3 dimBlock(h*w/num_thread, ch);
|
dim3 dimBlock(h*w/num_thread, ch);
|
||||||
normalize_kernel<<<dimBlock, num_thread, 0>>>(bgr, h*w, mean, stddev);
|
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);
|
|
||||||
}
|
|
||||||
@@ -1,41 +0,0 @@
|
|||||||
#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
|
|
||||||
);
|
|
||||||
|
|
||||||
}
|
|
||||||
}
|
|
||||||
@@ -1,6 +1,6 @@
|
|||||||
#include<iostream>
|
#include<iostream>
|
||||||
#include<algorithm>
|
#include<algorithm>
|
||||||
#include "tkdnn.h"
|
#include "tkDNN/tkdnn.h"
|
||||||
#include <stdlib.h> /* srand, rand */
|
#include <stdlib.h> /* srand, rand */
|
||||||
|
|
||||||
|
|
||||||
@@ -66,11 +66,11 @@ int main(int argc, char *argv[]) {
|
|||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
double min = *std::min_element(stats.begin(), stats.end())/BATCH_SIZE;
|
double min = *std::min_element(stats.begin(), stats.end()); ///BATCH_SIZE;
|
||||||
double max = *std::max_element(stats.begin(), stats.end())/BATCH_SIZE;
|
double max = *std::max_element(stats.begin(), stats.end()); ///BATCH_SIZE;
|
||||||
double mean =0;
|
double mean =0;
|
||||||
for(int i=0; i<stats.size(); i++) mean += stats[i]; mean /= stats.size();
|
for(int i=0; i<stats.size(); i++) mean += stats[i]; mean /= stats.size();
|
||||||
mean /=BATCH_SIZE;
|
//mean /=BATCH_SIZE;
|
||||||
|
|
||||||
std::cout<<"Min: "<<min<<" ms\n";
|
std::cout<<"Min: "<<min<<" ms\n";
|
||||||
std::cout<<"Max: "<<max<<" ms\n";
|
std::cout<<"Max: "<<max<<" ms\n";
|
||||||
|
|||||||
Reference in New Issue
Block a user