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

yolo3 berkeley ok

Browse Source
This commit is contained in:
Francesco Gatti
2019-02-18 15:37:39 +00:00
parent 2c63bf05be
commit 0d682136de
10 changed files with 278 additions and 459 deletions
Download Patch File Download Diff File Expand all files Collapse all files
CMakeLists.txt
+2 -5
View File
@@ -82,11 +82,8 @@ target_link_libraries(test_yolo3_berkeley tkDNN)
add_executable(test_rtinference tests/test_rtinference/rtinference.cpp) add_executable(test_rtinference tests/test_rtinference/rtinference.cpp)
target_link_libraries(test_rtinference tkDNN) target_link_libraries(test_rtinference tkDNN)
add_executable(detection demo/detection/detection.cpp) add_executable(yolo3_demo demo/demo/demo.cpp)
target_link_libraries(detection tkDNN) target_link_libraries(yolo3_demo tkDNN)
add_executable(live demo/live/live.cpp)
target_link_libraries(live tkDNN)
#install #install
#if (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT) #if (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT)
README.md
+11 -9
View File
@@ -1,11 +1,11 @@
# tkDNN # tkDNN
tkDNN is a Deep Neural Network library built with cuDNN primitives specifically thought to work on NVIDIA TK1 board.<br> tkDNN is a Deep Neural Network library built with cuDNN primitives specifically thought to work on NVIDIA TK1(and all successive) board.<br>
The main scope is to do high performance inference on already trained models. The main scope is to do high performance inference on already trained models.
this branch actually work on every NVIDIA GPU that support the dependencies: this branch actually work on every NVIDIA GPU that support the dependencies:
* CUDA 8 * CUDA 9
* CUDNN 6 * CUDNN 7.105
* TENSORRT 2 * TENSORRT 4.02
## Workflow ## Workflow
The recommended workflow follow these step: The recommended workflow follow these step:
@@ -32,18 +32,20 @@ Assumiung you have correctly builded the library these are the test ready to exe
* test_mnistRT: the mnist network hardcoded in using tensorRT apis (TENSORRT only) * test_mnistRT: the mnist network hardcoded in using tensorRT apis (TENSORRT only)
* test_yolo: YOLO detection network (CUDNN and TENSORRT) * test_yolo: YOLO detection network (CUDNN and TENSORRT)
* test_yolo_tiny: smaller version of YOLO (CUDNN and TENSRRT) * test_yolo_tiny: smaller version of YOLO (CUDNN and TENSRRT)
* test_yolo3_berkeley: our yolo3 version trained with BDD100K dateset
## Live detection ## yolo3 berkeley demo detection
For the live detection you need to precompile the tensorRT file by luncing the desidered network test, this is the recommended process: For the live detection you need to precompile the tensorRT file by luncing the desidered network test, this is the recommended process:
``` ```
export TKDNN_MODE=FP16 # set the half floating point optimization export TKDNN_MODE=FP16 # set the half floating point optimization
rm yolo.rt # be sure to delete(or move) old tensorRT files rm yolo3_berkeley.rt # be sure to delete(or move) old tensorRT files
./test_yolo # run the yolo test (is slow) ./test_yolo3_berkeley # run the yolo test (is slow)
# with f16 inference the result will be a bit incorrect # with f16 inference the result will be a bit incorrect
``` ```
this will genereate a yolo.rt file that can be used for live detection: this will genereate a yolo3_berkeley.rt file that can be used for live detection:
``` ```
./live yolo.rt 1 -s -t0.3 # launch detection on device 1 with 0.3 thresh ./demo # launch detection on a demo video
./demo /dev/video0 # launch detection on device 0
``` ```
demo/demo/demo.cpp
+75
View File
@@ -0,0 +1,75 @@
#include <iostream>
#include <signal.h>
#include <stdlib.h> /* srand, rand */
#include <unistd.h>
#include <mutex>
#include "utils.h"
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp>
#include "Yolo3Detection.h"
bool gRun;
void sig_handler(int signo) {
std::cout<<"request gateway stop\n";
gRun = false;
}
int main(int argc, char *argv[]) {
std::cout<<"detection\n";
signal(SIGINT, sig_handler);
Yolo3Detection yolo;
yolo.init("./");
gRun = true;
char *input = "../demo/yolo_test.mp4";
if(argc > 1)
input = argv[1];
cv::VideoCapture cap(input);
if(!cap.isOpened())
gRun = false;
else
std::cout<<"camera started\n";
cv::Mat frame;
cv::namedWindow("detection", cv::WINDOW_NORMAL);
cv::resizeWindow("detection", 544*1.2, 320*1.2);
while(gRun) {
cap >> frame;
if(!frame.data) {
continue;
}
yolo.update(frame);
// draw dets
for(int i=0; i<yolo.detected.size(); i++) {
tk::dnn::box b = yolo.detected[i];
int x0 = b.x;
int x1 = b.x + b.w;
int y0 = b.y;
int y1 = b.y + b.h;
int obj_class = b.cl;
float prob = b.prob;
std::cout<<obj_class<<" ("<<prob<<"): "<<x0<<" "<<y0<<" "<<x1<<" "<<y1<<"\n";
cv::rectangle(frame, cv::Point(x0, y0), cv::Point(x1, y1), yolo.colors[obj_class], 2);
}
cv::imshow("detection", frame);
cv::waitKey(1);
}
std::cout<<"detection end\n";
return 0;
}
demo/detection/detection.cpp
-249
View File
@@ -1,249 +0,0 @@
#include<iostream>
#include "tkdnn.h"
#include <stdlib.h> /* srand, rand */
#include <unistd.h>
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp>
const char *reg_bias = "../tests/yolo/layers/g31.bin";
int prob_sort(const void *pa, const void *pb) {
tk::dnn::box a = *(tk::dnn::box *)pa;
tk::dnn::box b = *(tk::dnn::box *)pb;
float diff = a.prob - b.prob;
if(diff < 0) return 1;
else if(diff > 0) return -1;
return 0;
}
cv::Mat GetSquareImage(const cv::Mat& img, int target_width) {
int width = img.cols, height = img.rows;
cv::Mat square = cv::Mat::zeros( target_width, target_width, img.type() );
int max_dim = ( width >= height ) ? width : height;
float scale = ( ( float ) target_width ) / max_dim;
cv::Rect roi;
if ( width >= height )
{
roi.width = target_width;
roi.x = 0;
roi.height = height * scale;
roi.y = ( target_width - roi.height ) / 2;
}
else
{
roi.y = 0;
roi.height = target_width;
roi.width = width * scale;
roi.x = ( target_width - roi.width ) / 2;
}
cv::resize( img, square( roi ), roi.size() );
return square;
}
//return inference time
double compute_image( cv::Mat imageORIG,
tk::dnn::NetworkRT *netRT, tk::dnn::RegionInterpret *rI,
dnnType *input, dnnType *output) {
//Resize with padding and convert to float
cv::Mat image = GetSquareImage(imageORIG, netRT->input_dim.w);
cv::Mat imageF;
image.convertTo(imageF, CV_32FC3, 1/255.0);
//split channels
cv::Mat bgr[3]; //destination array
cv::split(imageF,bgr);//split source
//write channels
int idx = 0;
memcpy((void*)&input[idx], (void*)bgr[2].data, imageF.rows*imageF.cols*sizeof(dnnType));
idx = imageF.rows*imageF.cols;
memcpy((void*)&input[idx], (void*)bgr[1].data, imageF.rows*imageF.cols*sizeof(dnnType));
idx *= 2;
memcpy((void*)&input[idx], (void*)bgr[0].data, imageF.rows*imageF.cols*sizeof(dnnType));
//DO INFERENCE
printCenteredTitle(" TENSORRT inference ", '=', 30);
TIMER_START
checkCuda( cudaMemcpyAsync(netRT->buffersRT[netRT->buf_input_idx], input,
netRT->input_dim.tot()*sizeof(float),
cudaMemcpyHostToDevice, netRT->stream));
netRT->enqueue();
checkCuda( cudaMemcpyAsync(output, netRT->buffersRT[netRT->buf_output_idx],
netRT->output_dim.tot()*sizeof(float),
cudaMemcpyDeviceToHost, netRT->stream));
cudaStreamSynchronize(netRT->stream);
TIMER_STOP
rI->interpretData(output, imageORIG.cols, imageORIG.rows);
return t_ns;
}
int print_usage() {
std::cout<<"usage: ./detection net.rt validation_list.txt"
<<" [-t <thresh>] [-s] [-i <iterations>]\n"
<<" -t: set thresh value\n -s: show images as compute\n"
<<" -i: images to compute\n\n"
<<"> validation_list.txt format: \n"
<<" path/to/image.jpg path/to/label.txt\n"
<<"> label.txt format: \n"
<<" <object-class> <x> <y> <width> <height>\n"
<<" x and y are the box center, "
<<"all values are relative to the image size\n\n";
return 1;
}
int main(int argc, char *argv[]) {
//params
char *tensor_path = NULL;
char *imageset_path = NULL;
float thresh = 0.3f;
bool show = false;
int iterations = INT_MAX;
//parse params
int c;
while ((c = getopt (argc, argv, "t:si:")) != -1) {
switch(c) {
case 't': thresh = atof(optarg); break;
case 's': show = true; break;
case 'i': iterations = atoi(optarg); break;
case '?':
return print_usage();
default: return print_usage();
}
}
if(argc - optind == 2) {
tensor_path = argv[optind];
imageset_path = argv[optind+1];
} else {
std::cout<<"not enough arguments.\n";
return print_usage();
}
//end parsing
if(!fileExist(tensor_path))
FatalError("unable to read serialRT file");
//convert network to tensorRT
tk::dnn::NetworkRT netRT(NULL, tensor_path);
tk::dnn::RegionInterpret rI(netRT.input_dim, netRT.output_dim, 80, 4, 5, thresh, reg_bias);
dnnType *input = new float[netRT.input_dim.tot()];
dnnType *output = new float[netRT.output_dim.tot()];
std::string line;
std::ifstream imageset(imageset_path);
if(!imageset.is_open())
FatalError("could not read imageset");
double mTime = 0;
float mAP = 0;
int processed_images;
for(processed_images=1;
processed_images-1 < iterations && getline(imageset, line);
processed_images++) {
std::string image_path = line.substr(0, line.find(" "));
std::string label_path = line.substr(line.find(" ")+1, line.size());
std::cout<<image_path<<"\n"<<label_path<<"\n";
//LOAD IMAGE
cv::Mat img = cv::imread(image_path.c_str(), CV_LOAD_IMAGE_COLOR);
if(!img.data)
FatalError("Could not open image");
std::cout<<"Image size: ("<<img.cols<<"x"<<img.rows<<")\n";
mTime += compute_image(img, &netRT, &rI, input, output);
std::ifstream labels(label_path.c_str());
if(!labels.is_open())
FatalError("could not read labels");
qsort(rI.res_boxes, rI.res_boxes_n, sizeof(tk::dnn::box), prob_sort);
for(int i=0; i<rI.res_boxes_n; i++) {
tk::dnn::box bx = rI.res_boxes[i];
std::cout<<" ("<<int(bx.prob*100)<<"%) "<<bx.cl
<<": "<<bx.x<<" "<<bx.y<<" "<<bx.w<<" "<<bx.h<<"\n";
cv::rectangle(img, cv::Point(bx.x - bx.w/2, bx.y - bx.h/2),
cv::Point(bx.x + bx.w/2, bx.y + bx.h/2),
cv::Scalar( 0, 0, 255), 2);
}
std::cout<<"GROUND TRUTH\n";
tk::dnn::box gt[256];
int gt_n = 0;
int cl;
float x, y, w, h;
while(labels>>cl) {
labels>>x>>y>>w>>h;
w *= img.cols; x *= img.cols;
h *= img.rows; y *= img.rows;
std::cout<<cl<<": "<<x<<" "<<y<<" "<<w<<" "<<h<<"\n";
gt[gt_n].x = x;
gt[gt_n].y = y;
gt[gt_n].w = w;
gt[gt_n].h = h;
gt[gt_n].cl = cl;
gt_n++;
cv::rectangle(img, cv::Point(x -w/2, y -h/2),
cv::Point(x +w/2, y +h/2),
cv::Scalar( 255, 0, 0), 2);
}
//AP calculation
float AP = 0;
for(int i=rI.res_boxes_n; i>=1; i--) { //for each detected evaluate sub group
int prec = 0;
for(int j=0; j<i; j++) { //for each detected in sub group
for(int z=0; z<gt_n; z++) { //control each ground truth
float iou = tk::dnn::RegionInterpret::box_iou(rI.res_boxes[j], gt[z]);
if(iou > 0.6f && rI.res_boxes[j].cl == gt[z].cl) {
prec++;
break;
}
}
}
AP += float(prec)/i;
}
AP = AP/gt_n;
std::cout<<"AP: "<<AP<<"\n";
mAP += AP;
std::cout<<"#### processed: "<<processed_images
<<", mAP: "<<mAP/processed_images<<"\n";
//show results
if(show) {
cv::namedWindow("result");
cv::imshow("result", img);
cv::waitKey(10);
}
}
//print results to file
processed_images -= 1;
std::ofstream res("results.txt", std::ios::app);
res<<"#### "<<tensor_path<<"\n";
res<<"processed images: "<<processed_images<<"\n";
res<<"mean inference time: "<<mTime/processed_images<<"\n";
res<<"mean AP: "<<mAP/processed_images<<"\n";
res<<"thesh used: "<<thresh<<"\n\n";
return 0;
}
demo/live/live.cpp
-196
View File
@@ -1,196 +0,0 @@
#include<iostream>
#include "tkdnn.h"
#include <stdlib.h> /* srand, rand */
#include <unistd.h>
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp>
#define VOC
#ifdef VOC
const char *reg_bias = "../tests/yolo_voc/layers/g31.bin";
#define CLASS 20
#else
const char *reg_bias = "../tests/yolo/layers/g31.bin";
#define CLASS 80
#endif
int prob_sort(const void *pa, const void *pb) {
tk::dnn::box a = *(tk::dnn::box *)pa;
tk::dnn::box b = *(tk::dnn::box *)pb;
float diff = a.prob - b.prob;
if(diff < 0) return 1;
else if(diff > 0) return -1;
return 0;
}
cv::Mat GetSquareImage(const cv::Mat& img, int target_width) {
int width = img.cols, height = img.rows;
cv::Mat square = cv::Mat::zeros( target_width, target_width, img.type() );
int max_dim = ( width >= height ) ? width : height;
float scale = ( ( float ) target_width ) / max_dim;
cv::Rect roi;
if ( width >= height )
{
roi.width = target_width;
roi.x = 0;
roi.height = height * scale;
roi.y = ( target_width - roi.height ) / 2;
}
else
{
roi.y = 0;
roi.height = target_width;
roi.width = width * scale;
roi.x = ( target_width - roi.width ) / 2;
}
cv::resize( img, square( roi ), roi.size() );
return square;
}
//return inference time
double compute_image( cv::Mat imageORIG,
tk::dnn::NetworkRT *netRT, tk::dnn::RegionInterpret *rI,
dnnType *input, dnnType *output) {
TIMER_START
//Resize with padding and convert to float
cv::Mat image = GetSquareImage(imageORIG, netRT->input_dim.w);
cv::Mat imageF;
image.convertTo(imageF, CV_32FC3, 1/255.0);
//split channels
cv::Mat bgr[3]; //destination array
cv::split(imageF,bgr);//split source
//write channels
int idx = 0;
memcpy((void*)&input[idx], (void*)bgr[2].data, imageF.rows*imageF.cols*sizeof(dnnType));
idx = imageF.rows*imageF.cols;
memcpy((void*)&input[idx], (void*)bgr[1].data, imageF.rows*imageF.cols*sizeof(dnnType));
idx *= 2;
memcpy((void*)&input[idx], (void*)bgr[0].data, imageF.rows*imageF.cols*sizeof(dnnType));
//DO INFERENCE
checkCuda( cudaMemcpyAsync(netRT->buffersRT[netRT->buf_input_idx], input,
netRT->input_dim.tot()*sizeof(float),
cudaMemcpyHostToDevice, netRT->stream));
netRT->enqueue();
checkCuda( cudaMemcpyAsync(output, netRT->buffersRT[netRT->buf_output_idx],
netRT->output_dim.tot()*sizeof(float),
cudaMemcpyDeviceToHost, netRT->stream));
cudaStreamSynchronize(netRT->stream);
rI->interpretData(output, imageORIG.cols, imageORIG.rows);
TIMER_STOP
return t_ns;
}
int print_usage() {
std::cout<<"usage: ./live net.rt input_uri\n";
return 1;
}
int main(int argc, char *argv[]) {
//params
char *tensor_path = NULL;
char *device = 0;
float thresh = 0.3f;
bool show = false;
//parse params
int c;
while ((c = getopt (argc, argv, "t:si:")) != -1) {
switch(c) {
case 't': thresh = atof(optarg); break;
case 's': show = true; break;
case '?':
return print_usage();
default: return print_usage();
}
}
if(argc - optind == 2) {
tensor_path = argv[optind];
device = argv[optind+1];
} else {
std::cout<<"not enough arguments.\n";
return print_usage();
}
//end parsing
std::cout<<"open video stream on device: "<<device<<"\n";
cv::VideoCapture cap(device);
/*
const char* pipe = "nvcamerasrc ! video/x-raw(memory:NVMM), width=(int)640, height=(int)480, format=(string)I420, framerate=(fraction)30/1 ! nvvidconv ! video/x-raw, format=(string)I420 ! videoconvert ! video/x-raw, format=(string)BGR ! appsink";
cv::VideoCapture cap(pipe);
*/
if(!cap.isOpened())
FatalError("unable to open video stream");
//cap.set(CV_CAP_PROP_BUFFERSIZE, 1); // process only last frame
if(!fileExist(tensor_path))
FatalError("unable to read serialRT file");
//convert network to tensorRT
tk::dnn::NetworkRT netRT(NULL, tensor_path);
tk::dnn::RegionInterpret rI(netRT.input_dim, netRT.output_dim, CLASS, 4, 5, thresh, reg_bias);
dnnType *input = new float[netRT.input_dim.tot()];
dnnType *output = new float[netRT.output_dim.tot()];
double mTime = 0;
int processed_images = 0;
for(;;) {
//LOAD IMAGE
cv::Mat img; //= cv::imread("../demo/live/test.jpeg", CV_LOAD_IMAGE_COLOR);
cap >> img;
if(!img.data)
FatalError("Could not open image");
std::cout<<"Image size: ("<<img.cols<<"x"<<img.rows<<")\n";
mTime += compute_image(img, &netRT, &rI, input, output);
qsort(rI.res_boxes, rI.res_boxes_n, sizeof(tk::dnn::box), prob_sort);
for(int i=0; i<rI.res_boxes_n; i++) {
tk::dnn::box bx = rI.res_boxes[i];
std::cout<<" ("<<int(bx.prob*100)<<"%) "<<bx.cl
<<": "<<bx.x<<" "<<bx.y<<" "<<bx.w<<" "<<bx.h<<"\n";
cv::rectangle(img, cv::Point(bx.x - bx.w/2, bx.y - bx.h/2),
cv::Point(bx.x + bx.w/2, bx.y + bx.h/2),
cv::Scalar( 0, 0, 255), 2);
}
//show results
if(show) {
cv::namedWindow("result");
cv::imshow("result", img);
cv::waitKey(1);
}
processed_images++;
std::cout<<"mean time per frames: "<<mTime/processed_images/1000<<" ms\n"<<"\n";
}
return 0;
}
demo/live/test.jpeg
BIN
View File
Binary file not shown.
Side by Side

Before

Width:  |  Height:  |  Size: 33 KiB

demo/yolo_test.mp4
BIN
View File
Binary file not shown.
include/Yolo3Detection.h
+53
View File
@@ -0,0 +1,53 @@
#include <iostream>
#include <signal.h>
#include <stdlib.h> /* srand, rand */
#include <unistd.h>
#include <mutex>
#include "utils.h"
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp>
#include <tkDNN/tkdnn.h>
/**
*
* @author Francesco Gatti
*/
class Yolo3Detection {
private:
tk::dnn::NetworkRT *netRT = nullptr;
tk::dnn::Yolo* yolo[3];
dnnType *input, *input_d;
int ndets = 0;
tk::dnn::Yolo::detection *dets = nullptr;
cv::Mat imageF;
cv::Mat bgr[3];
public:
static const int classes = 10;
static const int num = 3;
float thresh = 0.3;
cv::Scalar colors[classes];
// this is filled with results
std::vector<tk::dnn::box> detected;
Yolo3Detection() {}
virtual ~Yolo3Detection() {}
/**
* Method used for inizialize the class
*
* @return Success of the initialization
*/
bool init(std::string tensor_path);
void update(cv::Mat &frame);
};
src/Yolo3Detection.cpp
+128
View File
@@ -0,0 +1,128 @@
#include "Yolo3Detection.h"
bool Yolo3Detection::init(std::string tensor_folder) {
//const char *tensor_path = "../data/yolo3/yolo3_berkeley.rt";
// class colors precompute
for(int c=0; c<classes; c++) {
int cc = c+1;
double d = 1.0*( (cc%16)/8 );
double r = 1.0*( (cc%8)/4 ) + (0.5*d);
double g = 1.0*( (cc%4)/2 ) + (0.5*d);
double b = 1.0*( (cc%2)/1 ) + (0.5*d);
if(r > 1) r = 1;
if(g > 1) g = 1;
if(b > 1) b = 1;
//std::cout<<r<<" "<<g<<" "<<b<<"\n";
colors[c] = cv::Scalar(int(255.0*b), int(255.0*g), int(255.0*r));
}
//convert network to tensorRT
std::cout<<(tensor_folder + "/yolo3_berkeley.rt").c_str()<<"\n";
netRT = new tk::dnn::NetworkRT(NULL, (tensor_folder + "/yolo3_berkeley.rt").c_str() );
yolo[0] = new tk::dnn::Yolo(nullptr, classes, num, (tensor_folder + "/yolo3_0.bin").c_str() ); // yolo without input and bias
yolo[0]->input_dim = yolo[0]->output_dim = tk::dnn::dataDim_t(1, 45, 10, 17);
yolo[1] = new tk::dnn::Yolo(nullptr, classes, num, (tensor_folder + "/yolo3_1.bin").c_str() ); // yolo without input and bias
yolo[1]->input_dim = yolo[1]->output_dim = tk::dnn::dataDim_t(1, 45, 20, 34);
yolo[2] = new tk::dnn::Yolo(nullptr, classes, num, (tensor_folder + "/yolo3_2.bin").c_str() ); // yolo without input and bias
yolo[2]->input_dim = yolo[2]->output_dim = tk::dnn::dataDim_t(1, 45, 40, 68);
dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes);
checkCuda(cudaMallocHost(&input, sizeof(dnnType)*netRT->input_dim.tot()));
checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot()));
return true;
}
void Yolo3Detection::update(cv::Mat &imageORIG) {
if(!imageORIG.data) {
std::cout<<"YOLO: NO IMAGE DATA\n";
return;
}
resize(imageORIG, imageORIG, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
imageORIG.convertTo(imageF, CV_32FC3, 1/255.0);
//split channels
cv::split(imageF,bgr);//split source
//write channels
int idx = 0;
memcpy((void*)&input[idx], (void*)bgr[2].data, imageF.rows*imageF.cols*sizeof(dnnType));
idx = imageF.rows*imageF.cols;
memcpy((void*)&input[idx], (void*)bgr[1].data, imageF.rows*imageF.cols*sizeof(dnnType));
idx *= 2;
memcpy((void*)&input[idx], (void*)bgr[0].data, imageF.rows*imageF.cols*sizeof(dnnType));
//DO INFERENCE
dnnType *rt_out[3];
tk::dnn::dataDim_t dim = netRT->input_dim;
checkCuda(cudaMemcpyAsync(input_d, input, dim.tot()*sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
printCenteredTitle(" TENSORRT inference ", '=', 30); {
dim.print();
TIMER_START
netRT->infer(dim, input_d);
TIMER_STOP
dim.print();
}
TIMER_START
// compute dets
ndets = 0;
for(int i=0; i<3; i++) {
rt_out[i] = (dnnType*)netRT->buffersRT[i+1];
yolo[i]->dstData = rt_out[i];
yolo[i]->computeDetections(dets, ndets, netRT->input_dim.w, netRT->input_dim.h, netRT->input_dim.w, netRT->input_dim.h, thresh);
}
tk::dnn::Yolo::mergeDetections(dets, ndets, classes);
TIMER_STOP
float xRatio = float(imageORIG.cols) / float(netRT->input_dim.w);
float yRatio = float(imageORIG.rows) / float(netRT->input_dim.h);
// fill detected
detected.clear();
for(int j=0; j<ndets; j++) {
tk::dnn::Yolo::box b = dets[j].bbox;
int x0 = (b.x-b.w/2.);
int x1 = (b.x+b.w/2.);
int y0 = (b.y-b.h/2.);
int y1 = (b.y+b.h/2.);
int obj_class = -1;
float prob = 0;
for(int c=0; c<classes; c++) {
if(dets[j].prob[c] >= thresh) {
obj_class = c;
prob = dets[j].prob[c];
}
}
if(obj_class >= 0) {
//std::cout<<obj_class<<" ("<<prob<<"): "<<x0<<" "<<y0<<" "<<x1<<" "<<y1<<"\n";
//cv::rectangle(image, cv::Point(x0, y0), cv::Point(x1, y1), colors[obj_class], 2);
// convert to image coords
x0 = xRatio*x0;
x1 = xRatio*x1;
y0 = yRatio*y0;
y1 = yRatio*y1;
tk::dnn::box res;
res.cl = obj_class;
res.prob = prob;
res.x = x0;
res.y = y0;
res.w = x1 - x0;
res.h = y1 - y0;
detected.push_back(res);
}
}
}
tests/yolo3_berkeley/yolo3_berkeley.cpp
+9
View File
@@ -369,5 +369,14 @@ int main() {
std::cout<<"TRT vs correct"; checkResult(odim, rt_out[i], out); std::cout<<"TRT vs correct"; checkResult(odim, rt_out[i], out);
std::cout<<"CUDNN vs TRT "; checkResult(odim, cudnn_out[i], rt_out[i]); std::cout<<"CUDNN vs TRT "; checkResult(odim, cudnn_out[i], rt_out[i]);
} }
std::cout<<"copyng layer config to this folder\n";
std::string cmd;
cmd = "cp " + std::string(g82_bin) + " yolo3_0.bin";
std::cout<<cmd<<"\n"; system(cmd.c_str());
cmd = "cp " + std::string(g94_bin) + " yolo3_1.bin";
std::cout<<cmd<<"\n"; system(cmd.c_str());
cmd = "cp " + std::string(g106_bin) + " yolo3_2.bin";
std::cout<<cmd<<"\n"; system(cmd.c_str());
return 0; return 0;
} }