Merge branch 'cnet' of https://github.com/ceccocats/tkDNN into cnet
This commit is contained in:
+18
-1
@@ -6,14 +6,31 @@
|
|||||||
#include <thrust/gather.h>
|
#include <thrust/gather.h>
|
||||||
#include <thrust/copy.h>
|
#include <thrust/copy.h>
|
||||||
|
|
||||||
|
#include <opencv2/core/core.hpp>
|
||||||
|
#include <opencv2/highgui/highgui.hpp>
|
||||||
|
#include <opencv2/imgproc/imgproc.hpp>
|
||||||
|
#include "opencv2/opencv.hpp"
|
||||||
|
|
||||||
#include "tkdnn.h"
|
#include "tkdnn.h"
|
||||||
|
|
||||||
|
struct threshold : public thrust::binary_function<float,float,float>
|
||||||
|
{
|
||||||
|
__host__ __device__
|
||||||
|
float operator()(float x, float y) {
|
||||||
|
double toll = 1e-6;
|
||||||
|
if(fabsf(x-y)>toll)
|
||||||
|
return 0.0f;
|
||||||
|
else
|
||||||
|
return x;
|
||||||
|
}
|
||||||
|
};
|
||||||
|
|
||||||
void sort(dnnType *src_begin, dnnType *src_end, int *idsrc);
|
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 subtractWithThreshold(dnnType *src_begin, dnnType *src_end, dnnType *src2_begin, dnnType *src_out);
|
void normalize(float *bgr, const int ch, const int h, const int w, const float *mean, const float *stddev);
|
||||||
|
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);
|
||||||
void topKxyAddOffset(int * ids_begin, const int K, const int size, int *intxs_begin, int *intys_begin,
|
void topKxyAddOffset(int * ids_begin, const int K, const int size, int *intxs_begin, int *intys_begin,
|
||||||
float *xs_begin, float *ys_begin, dnnType *src_begin, float *src_out, int *ids_out);
|
float *xs_begin, float *ys_begin, dnnType *src_begin, float *src_out, int *ids_out);
|
||||||
|
|||||||
@@ -15,6 +15,7 @@
|
|||||||
#include <opencv2/core/core.hpp>
|
#include <opencv2/core/core.hpp>
|
||||||
#include <opencv2/highgui/highgui.hpp>
|
#include <opencv2/highgui/highgui.hpp>
|
||||||
#include <opencv2/imgproc/imgproc.hpp>
|
#include <opencv2/imgproc/imgproc.hpp>
|
||||||
|
#include "opencv2/opencv.hpp"
|
||||||
|
|
||||||
#include "tkdnn.h"
|
#include "tkdnn.h"
|
||||||
#include "sorting.h"
|
#include "sorting.h"
|
||||||
@@ -29,20 +30,24 @@ class CenternetDetection {
|
|||||||
|
|
||||||
private:
|
private:
|
||||||
tk::dnn::NetworkRT *netRT = nullptr;
|
tk::dnn::NetworkRT *netRT = nullptr;
|
||||||
dnnType *input_h, *input, *input_d;
|
dnnType *input_d;
|
||||||
|
|
||||||
int ndets = 0;
|
int ndets = 0;
|
||||||
// tk::dnn::Yolo::detection *dets = nullptr;
|
// tk::dnn::Yolo::detection *dets = nullptr;
|
||||||
|
|
||||||
cv::Mat imageF;
|
cv::Mat imageF;
|
||||||
cv::Mat bgr[3];
|
cv::cuda::GpuMat imageF1_d, imageF2_d;
|
||||||
|
cv::cuda::GpuMat bgr[3];
|
||||||
|
// std::vector< cv::cuda::GpuMat > bgr;
|
||||||
|
|
||||||
// variable to test cnet on dog pictures
|
// variable to test cnet on dog pictures
|
||||||
tk::dnn::dataDim_t dim;
|
tk::dnn::dataDim_t dim;
|
||||||
tk::dnn::dataDim_t dim2;
|
tk::dnn::dataDim_t dim2;
|
||||||
cv::Size sz;
|
cv::Size sz, sz_old;
|
||||||
const char *input_bin = "../tests/resnet101_cnet/debug/input.bin";
|
const char *input_bin = "../tests/resnet101_cnet/debug/input.bin";
|
||||||
|
|
||||||
|
cv::cuda::Stream stream;
|
||||||
|
struct threshold op;
|
||||||
// pre-process
|
// pre-process
|
||||||
tk::dnn::dataDim_t dim_hm;
|
tk::dnn::dataDim_t dim_hm;
|
||||||
tk::dnn::dataDim_t dim_wh;
|
tk::dnn::dataDim_t dim_wh;
|
||||||
@@ -66,10 +71,15 @@ class CenternetDetection {
|
|||||||
|
|
||||||
float *target_coords;
|
float *target_coords;
|
||||||
|
|
||||||
cv::Vec<float, 3> mean;
|
float *mean_d;
|
||||||
cv::Vec<float, 3> stddev;
|
float *stddev_d;
|
||||||
|
|
||||||
|
float *d_ptrs;
|
||||||
|
|
||||||
cv::Mat src;
|
cv::Mat src;
|
||||||
cv::Mat dst;
|
cv::Mat dst;
|
||||||
|
cv::Mat dst2;
|
||||||
|
cv::Mat trans, trans2;
|
||||||
//processing
|
//processing
|
||||||
float toll = 0.000001;
|
float toll = 0.000001;
|
||||||
int K = 100;
|
int K = 100;
|
||||||
@@ -108,7 +118,6 @@ class CenternetDetection {
|
|||||||
* @return Success of the initialization
|
* @return Success of the initialization
|
||||||
*/
|
*/
|
||||||
bool init(std::string tensor_path);
|
bool init(std::string tensor_path);
|
||||||
void testdog();
|
|
||||||
cv::Mat draw(cv::Mat &frame);
|
cv::Mat draw(cv::Mat &frame);
|
||||||
void update(cv::Mat &frame);
|
void update(cv::Mat &frame);
|
||||||
|
|
||||||
|
|||||||
+101
-144
@@ -37,10 +37,11 @@ bool CenternetDetection::init(std::string tensor_path) {
|
|||||||
coco_class_name = std::vector<std::string>(coco_class_name_, std::end( coco_class_name_ ));
|
coco_class_name = std::vector<std::string>(coco_class_name_, std::end( coco_class_name_ ));
|
||||||
src = cv::Mat(cv::Size(2,3), CV_32F);
|
src = cv::Mat(cv::Size(2,3), CV_32F);
|
||||||
dst = cv::Mat(cv::Size(2,3), CV_32F);
|
dst = cv::Mat(cv::Size(2,3), CV_32F);
|
||||||
|
dst2 = cv::Mat(cv::Size(2,3), CV_32F);
|
||||||
|
trans = cv::Mat(cv::Size(3,2), CV_32F);
|
||||||
|
trans2 = cv::Mat(cv::Size(3,2), CV_32F);
|
||||||
// dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes);
|
// dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes);
|
||||||
|
|
||||||
checkCuda(cudaMallocHost(&input_h, sizeof(dnnType)*netRT->input_dim.tot()));
|
|
||||||
checkCuda(cudaMallocHost(&input, sizeof(dnnType)*netRT->input_dim.tot()));
|
|
||||||
checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot()));
|
checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot()));
|
||||||
|
|
||||||
// dim_hm = tk::dnn::dataDim_t(1, 80, 56, 56, 1);
|
// dim_hm = tk::dnn::dataDim_t(1, 80, 56, 56, 1);
|
||||||
@@ -94,61 +95,31 @@ bool CenternetDetection::init(std::string tensor_path) {
|
|||||||
|
|
||||||
checkCuda( cudaMallocHost(&target_coords, 4 * K *sizeof(float)) );
|
checkCuda( cudaMallocHost(&target_coords, 4 * K *sizeof(float)) );
|
||||||
|
|
||||||
mean << 0.408, 0.447, 0.47;
|
checkCuda( cudaMalloc(&mean_d, 3 * sizeof(float)) );
|
||||||
stddev << 0.289, 0.274, 0.278;
|
checkCuda( cudaMalloc(&stddev_d, 3 * sizeof(float)) );
|
||||||
|
float mean[3] = {0.408, 0.447, 0.47};
|
||||||
|
float stddev[3] = {0.289, 0.274, 0.278};
|
||||||
|
|
||||||
|
checkCuda(cudaMemcpy(mean_d, mean, 3*sizeof(float), cudaMemcpyHostToDevice));
|
||||||
|
checkCuda(cudaMemcpy(stddev_d, stddev, 3*sizeof(float), cudaMemcpyHostToDevice));
|
||||||
|
|
||||||
|
checkCuda( cudaMalloc(&d_ptrs, dim.c * dim.h*dim.w * sizeof(float)) );
|
||||||
|
// mean << 0.408, 0.447, 0.47;
|
||||||
|
// stddev << 0.289, 0.274, 0.278;
|
||||||
|
|
||||||
// Alloc array used in the kernel
|
// Alloc array used in the kernel
|
||||||
checkCuda( cudaMalloc(&src_out, K *sizeof(float)) );
|
checkCuda( cudaMalloc(&src_out, K *sizeof(float)) );
|
||||||
checkCuda( cudaMalloc(&ids_out, K *sizeof(int)) );
|
checkCuda( cudaMalloc(&ids_out, K *sizeof(int)) );
|
||||||
// checkCuda( cudaFree(src_out) );
|
// checkCuda( cudaFree(src_out) );
|
||||||
// checkCuda( cudaFree(ids_out) );
|
// checkCuda( cudaFree(ids_out) );
|
||||||
|
dst2.at<float>(0,0)=width * 0.5;
|
||||||
}
|
dst2.at<float>(0,1)=width * 0.5;
|
||||||
|
dst2.at<float>(1,0)=width * 0.5;
|
||||||
void CenternetDetection::testdog() {
|
dst2.at<float>(1,1)=width * 0.5 + width * -0.5;
|
||||||
|
|
||||||
readBinaryFile(input_bin, dim.tot(), &input_h, &input_d);
|
dst2.at<float>(2,0)=dst2.at<float>(1,0) + (-dst2.at<float>(0,1)+dst2.at<float>(1,1) );
|
||||||
|
dst2.at<float>(2,1)=dst2.at<float>(1,1) + (dst2.at<float>(0,0)-dst2.at<float>(1,0) );
|
||||||
// -------- transofrm compose
|
|
||||||
cv::Mat imageORIG = cv::imread("../../dog.jpg");
|
|
||||||
imageORIG.convertTo(imageF, CV_32FC3, 1/255.0);
|
|
||||||
sz = imageF.size();
|
|
||||||
std::cout<<"image: "<<sz.width<<", "<<sz.height<<std::endl;
|
|
||||||
resize(imageF, imageF, cv::Size(512, 512));
|
|
||||||
const int cropSize = 512;
|
|
||||||
const int offsetW = (imageF.cols - cropSize) / 2;
|
|
||||||
const int offsetH = (imageF.rows - cropSize) / 2;
|
|
||||||
const cv::Rect roi(offsetW, offsetH, cropSize, cropSize);
|
|
||||||
imageF = imageF(roi).clone();
|
|
||||||
std::cout << "Cropped image dimension: " << imageF.cols << " X " << imageF.rows << std::endl;
|
|
||||||
|
|
||||||
mean << 0.485, 0.456, 0.406;
|
|
||||||
stddev << 0.229, 0.224, 0.225;
|
|
||||||
sz = imageF.size();
|
|
||||||
// std::cout<<"size: "<<sz.height<<" "<<sz.width<<" - "<<std::endl;
|
|
||||||
// std::cout<<"mean: "<<mean<<", std: "<<stddev<<std::endl;
|
|
||||||
cv::add(imageF, -mean, imageF);
|
|
||||||
cv::divide(imageF, stddev, imageF);
|
|
||||||
//split channels
|
|
||||||
cv::split(imageF,bgr);//split source
|
|
||||||
dim2 = dim;
|
|
||||||
//write channels
|
|
||||||
for(int i=0; i<dim2.c; i++) {
|
|
||||||
int idx = i*imageF.rows*imageF.cols;
|
|
||||||
int ch = dim2.c-1 -i;
|
|
||||||
memcpy((void*)&input[idx], (void*)bgr[ch].data, imageF.rows*imageF.cols*sizeof(dnnType));
|
|
||||||
}
|
|
||||||
|
|
||||||
checkCuda(cudaMemcpyAsync(input_d, input, dim2.tot()*sizeof(dnnType), cudaMemcpyHostToDevice));
|
|
||||||
|
|
||||||
printCenteredTitle(" TENSORRT inference ", '=', 30); {
|
|
||||||
dim2.print();
|
|
||||||
TIMER_START
|
|
||||||
netRT->infer(dim2, input_d);
|
|
||||||
TIMER_STOP
|
|
||||||
dim2.print();
|
|
||||||
}
|
|
||||||
// checkResult(dim2.tot(), input_h, input);
|
|
||||||
}
|
}
|
||||||
|
|
||||||
cv::Mat CenternetDetection::draw(cv::Mat &imageORIG) {
|
cv::Mat CenternetDetection::draw(cv::Mat &imageORIG) {
|
||||||
@@ -196,7 +167,7 @@ cv::Mat CenternetDetection::draw(cv::Mat &imageORIG) {
|
|||||||
void CenternetDetection::update(cv::Mat &imageORIG) {
|
void CenternetDetection::update(cv::Mat &imageORIG) {
|
||||||
|
|
||||||
if(!imageORIG.data) {
|
if(!imageORIG.data) {
|
||||||
std::cout<<"YOLO: NO IMAGE DATA\n";
|
std::cout<<"CENTERNET: NO IMAGE DATA\n";
|
||||||
return;
|
return;
|
||||||
}
|
}
|
||||||
TIMER_START
|
TIMER_START
|
||||||
@@ -211,82 +182,91 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
|
|||||||
float scale = 1.0;
|
float scale = 1.0;
|
||||||
float new_height = sz.height * scale;
|
float new_height = sz.height * scale;
|
||||||
float new_width = sz.width * scale;
|
float new_width = sz.width * scale;
|
||||||
float c[] = {new_width / 2.0, new_height /2.0};
|
if(sz.height != sz_old.height && sz.width != sz_old.width){
|
||||||
float s[2];
|
float c[] = {new_width / 2.0, new_height /2.0};
|
||||||
|
float s[2];
|
||||||
|
|
||||||
if(sz.width > sz.height){
|
if(sz.width > sz.height){
|
||||||
s[0] = sz.width * 1.0;
|
s[0] = sz.width * 1.0;
|
||||||
s[1] = sz.width * 1.0;
|
s[1] = sz.width * 1.0;
|
||||||
|
}
|
||||||
|
else{
|
||||||
|
s[0] = sz.height * 1.0;
|
||||||
|
s[1] = sz.height * 1.0;
|
||||||
|
}
|
||||||
|
|
||||||
|
// ----------- get_affine_transform
|
||||||
|
// rot_rad = pi * 0 / 100 --> 0
|
||||||
|
|
||||||
|
src.at<float>(0,0)=c[0];
|
||||||
|
src.at<float>(0,1)=c[1];
|
||||||
|
src.at<float>(1,0)=c[0];
|
||||||
|
src.at<float>(1,1)=c[1] + s[0] * -0.5;
|
||||||
|
dst.at<float>(0,0)=inp_width * 0.5;
|
||||||
|
dst.at<float>(0,1)=inp_height * 0.5;
|
||||||
|
dst.at<float>(1,0)=inp_width * 0.5;
|
||||||
|
dst.at<float>(1,1)=inp_height * 0.5 + inp_width * -0.5;
|
||||||
|
|
||||||
|
src.at<float>(2,0)=src.at<float>(1,0) + (-src.at<float>(0,1)+src.at<float>(1,1) );
|
||||||
|
src.at<float>(2,1)=src.at<float>(1,1) + (src.at<float>(0,0)-src.at<float>(1,0) );
|
||||||
|
dst.at<float>(2,0)=dst.at<float>(1,0) + (-dst.at<float>(0,1)+dst.at<float>(1,1) );
|
||||||
|
dst.at<float>(2,1)=dst.at<float>(1,1) + (dst.at<float>(0,0)-dst.at<float>(1,0) );
|
||||||
|
|
||||||
|
trans = cv::getAffineTransform( src, dst );
|
||||||
|
end_t = std::chrono::steady_clock::now();
|
||||||
|
std::cout << " TIME gett affine trans: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
|
step_t = end_t;
|
||||||
|
|
||||||
|
trans2 = cv::getAffineTransform( dst2, src );
|
||||||
|
|
||||||
|
end_t = std::chrono::steady_clock::now();
|
||||||
|
std::cout << " TIME getAffineTrans 2: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
|
step_t = end_t;
|
||||||
}
|
}
|
||||||
else{
|
sz_old = sz;
|
||||||
s[0] = sz.height * 1.0;
|
cv::cuda::GpuMat im_Orig;
|
||||||
s[1] = sz.height * 1.0;
|
im_Orig = cv::cuda::GpuMat(imageORIG);
|
||||||
}
|
cv::cuda::resize (im_Orig, imageF1_d, cv::Size(new_width, new_height));
|
||||||
|
checkCuda( cudaDeviceSynchronize() );
|
||||||
// ----------- get_affine_transform
|
|
||||||
// rot_rad = pi * 0 / 100 --> 0
|
|
||||||
|
|
||||||
src.at<float>(0,0)=c[0];
|
|
||||||
src.at<float>(0,1)=c[1];
|
|
||||||
src.at<float>(1,0)=c[0];
|
|
||||||
src.at<float>(1,1)=c[1] + s[0] * -0.5;
|
|
||||||
dst.at<float>(0,0)=inp_width * 0.5;
|
|
||||||
dst.at<float>(0,1)=inp_height * 0.5;
|
|
||||||
dst.at<float>(1,0)=inp_width * 0.5;
|
|
||||||
dst.at<float>(1,1)=inp_height * 0.5 + inp_width * -0.5;
|
|
||||||
|
|
||||||
src.at<float>(2,0)=src.at<float>(1,0) + (-src.at<float>(0,1)+src.at<float>(1,1) );
|
sz = imageF1_d.size();
|
||||||
src.at<float>(2,1)=src.at<float>(1,1) + (src.at<float>(0,0)-src.at<float>(1,0) );
|
|
||||||
dst.at<float>(2,0)=dst.at<float>(1,0) + (-dst.at<float>(0,1)+dst.at<float>(1,1) );
|
|
||||||
dst.at<float>(2,1)=dst.at<float>(1,1) + (dst.at<float>(0,0)-dst.at<float>(1,0) );
|
|
||||||
|
|
||||||
cv::Mat trans = cv::getAffineTransform( src, dst );
|
|
||||||
end_t = std::chrono::steady_clock::now();
|
|
||||||
std::cout << " TIME gett affine trans: " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
|
|
||||||
step_t = end_t;
|
|
||||||
|
|
||||||
|
|
||||||
resize(imageORIG, imageF, cv::Size(new_width, new_height));
|
|
||||||
sz = imageF.size();
|
|
||||||
std::cout<<"size: "<<sz.height<<" "<<sz.width<<" - "<<std::endl;
|
std::cout<<"size: "<<sz.height<<" "<<sz.width<<" - "<<std::endl;
|
||||||
end_t = std::chrono::steady_clock::now();
|
end_t = std::chrono::steady_clock::now();
|
||||||
std::cout << " TIME resize: " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
|
std::cout << " TIME resize: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
step_t = end_t;
|
step_t = end_t;
|
||||||
|
|
||||||
|
cv::cuda::warpAffine(imageF1_d, imageF2_d, trans, cv::Size(inp_width, inp_height), cv::INTER_LINEAR );
|
||||||
cv::warpAffine(imageF, imageF, trans, cv::Size(inp_width, inp_height), cv::INTER_LINEAR );
|
checkCuda( cudaDeviceSynchronize() );
|
||||||
end_t = std::chrono::steady_clock::now();
|
end_t = std::chrono::steady_clock::now();
|
||||||
std::cout << " TIME warpAffine: " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
|
std::cout << " TIME warpAffine: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
step_t = end_t;
|
step_t = end_t;
|
||||||
|
|
||||||
|
imageF2_d.convertTo(imageF1_d, CV_32FC3, 1/255.0);
|
||||||
sz = imageF.size();
|
checkCuda( cudaDeviceSynchronize() );
|
||||||
std::cout<<"size: "<<sz.height<<" "<<sz.width<<" - "<<std::endl;
|
|
||||||
imageF.convertTo(imageF, CV_32FC3, 1/255.0);
|
|
||||||
|
|
||||||
end_t = std::chrono::steady_clock::now();
|
end_t = std::chrono::steady_clock::now();
|
||||||
std::cout << " TIME convert: " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
|
std::cout << " TIME convert: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
step_t = end_t;
|
step_t = end_t;
|
||||||
|
|
||||||
dim2 = dim;
|
dim2 = dim;
|
||||||
|
cv::cuda::split(imageF1_d,bgr);//split source
|
||||||
//split channels
|
end_t = std::chrono::steady_clock::now();
|
||||||
cv::split(imageF,bgr);//split source
|
std::cout << " TIME split: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
|
step_t = end_t;
|
||||||
for(int i=0; i<3; i++){
|
|
||||||
bgr[i] = bgr[i] - mean[i];
|
|
||||||
bgr[i] = bgr[i] / stddev[i];
|
|
||||||
}
|
|
||||||
|
|
||||||
//write channels
|
for(int i=0; i<dim.c; i++)
|
||||||
for(int i=0; i<dim2.c; i++) {
|
checkCuda( cudaMemcpy(d_ptrs + i*dim.h * dim.w, (float*)bgr[i].data, dim.h * dim.w * sizeof(float), cudaMemcpyDeviceToDevice) );
|
||||||
int idx = i*imageF.rows*imageF.cols;
|
|
||||||
int ch = dim2.c-3 +i;
|
normalize(d_ptrs, dim.c, dim.h, dim.w, mean_d, stddev_d);
|
||||||
// std::cout<<"i: "<<i<<", idx: "<<idx<<", ch: "<<ch<<std::endl;
|
|
||||||
memcpy((void*)&input[idx], (void*)bgr[ch].data, imageF.rows*imageF.cols*sizeof(dnnType));
|
end_t = std::chrono::steady_clock::now();
|
||||||
}
|
std::cout << " TIME normalize: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
|
step_t = end_t;
|
||||||
|
|
||||||
checkCuda(cudaMemcpyAsync(input_d, input, dim2.tot()*sizeof(dnnType), cudaMemcpyHostToDevice));
|
checkCuda(cudaMemcpy(input_d, d_ptrs, dim2.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice));
|
||||||
|
|
||||||
|
end_t = std::chrono::steady_clock::now();
|
||||||
|
std::cout << " TIME Memcpy to input_d: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
|
step_t = end_t;
|
||||||
|
|
||||||
printCenteredTitle(" TENSORRT inference ", '=', 30); {
|
printCenteredTitle(" TENSORRT inference ", '=', 30); {
|
||||||
dim2.print();
|
dim2.print();
|
||||||
@@ -295,7 +275,6 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
|
|||||||
TIMER_STOP
|
TIMER_STOP
|
||||||
dim2.print();
|
dim2.print();
|
||||||
}
|
}
|
||||||
// checkResult(dim2.tot(), input_h, input);
|
|
||||||
step_t = std::chrono::steady_clock::now();
|
step_t = std::chrono::steady_clock::now();
|
||||||
|
|
||||||
// ------------------------------------ process --------------------------------------------
|
// ------------------------------------ process --------------------------------------------
|
||||||
@@ -307,10 +286,10 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
|
|||||||
activationSIGMOIDForward(rt_out[0], rt_out[0], dim_hm.tot());
|
activationSIGMOIDForward(rt_out[0], rt_out[0], dim_hm.tot());
|
||||||
checkCuda( cudaDeviceSynchronize() );
|
checkCuda( cudaDeviceSynchronize() );
|
||||||
|
|
||||||
subtractWithThreshold(rt_out[0], rt_out[0] + dim_hm.tot(), rt_out[1], rt_out[0]);
|
subtractWithThreshold(rt_out[0], rt_out[0] + dim_hm.tot(), rt_out[1], rt_out[0], op);
|
||||||
|
|
||||||
end_t = std::chrono::steady_clock::now();
|
end_t = std::chrono::steady_clock::now();
|
||||||
std::cout << " TIME threshold: " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
|
std::cout << " TIME threshold: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
step_t = end_t;
|
step_t = end_t;
|
||||||
// ----------- nms end
|
// ----------- nms end
|
||||||
// ----------- topk
|
// ----------- topk
|
||||||
@@ -327,7 +306,7 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
|
|||||||
ids_d);
|
ids_d);
|
||||||
checkCuda( cudaDeviceSynchronize() );
|
checkCuda( cudaDeviceSynchronize() );
|
||||||
end_t = std::chrono::steady_clock::now();
|
end_t = std::chrono::steady_clock::now();
|
||||||
std::cout << " TIME sort: " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
|
std::cout << " TIME sort: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
step_t = end_t;
|
step_t = end_t;
|
||||||
|
|
||||||
topk(rt_out[0], ids_d, K, scores_d,
|
topk(rt_out[0], ids_d, K, scores_d,
|
||||||
@@ -335,14 +314,14 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
|
|||||||
checkCuda( cudaDeviceSynchronize() );
|
checkCuda( cudaDeviceSynchronize() );
|
||||||
|
|
||||||
end_t = std::chrono::steady_clock::now();
|
end_t = std::chrono::steady_clock::now();
|
||||||
std::cout << " TIME topk: " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
|
std::cout << " TIME topk: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
step_t = end_t;
|
step_t = end_t;
|
||||||
|
|
||||||
checkCuda( cudaMemcpy(scores, scores_d, K *sizeof(float), cudaMemcpyDeviceToHost) );
|
checkCuda( cudaMemcpy(scores, scores_d, K *sizeof(float), cudaMemcpyDeviceToHost) );
|
||||||
|
|
||||||
topKxyclasses(topk_inds_d, topk_inds_d+K, K, width, dim_hm.w*dim_hm.h, clses_d, inttopk_xs_d, inttopk_ys_d);
|
topKxyclasses(topk_inds_d, topk_inds_d+K, K, width, dim_hm.w*dim_hm.h, clses_d, inttopk_xs_d, inttopk_ys_d);
|
||||||
end_t = std::chrono::steady_clock::now();
|
end_t = std::chrono::steady_clock::now();
|
||||||
std::cout << " TIME topk x y clses 2: " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
|
std::cout << " TIME topk x y clses 2: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
step_t = end_t;
|
step_t = end_t;
|
||||||
|
|
||||||
checkCuda( cudaMemcpy(topk_xs_d, (float *)inttopk_xs_d, K*sizeof(float), cudaMemcpyDeviceToDevice) );
|
checkCuda( cudaMemcpy(topk_xs_d, (float *)inttopk_xs_d, K*sizeof(float), cudaMemcpyDeviceToDevice) );
|
||||||
@@ -356,7 +335,7 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
|
|||||||
// checkCuda( cudaDeviceSynchronize() );
|
// checkCuda( cudaDeviceSynchronize() );
|
||||||
|
|
||||||
end_t = std::chrono::steady_clock::now();
|
end_t = std::chrono::steady_clock::now();
|
||||||
std::cout << " TIME add offset: " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
|
std::cout << " TIME add offset: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
step_t = end_t;
|
step_t = end_t;
|
||||||
|
|
||||||
bboxes(topk_inds_d, K, dim_wh.h*dim_wh.w, topk_xs_d, topk_ys_d, rt_out[2], bbx0_d, bbx1_d, bby0_d, bby1_d, src_out, ids_out);
|
bboxes(topk_inds_d, K, dim_wh.h*dim_wh.w, topk_xs_d, topk_ys_d, rt_out[2], bbx0_d, bbx1_d, bby0_d, bby1_d, src_out, ids_out);
|
||||||
@@ -368,35 +347,13 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
|
|||||||
checkCuda( cudaMemcpy(bby1, bby1_d, K * sizeof(float), cudaMemcpyDeviceToHost) );
|
checkCuda( cudaMemcpy(bby1, bby1_d, K * sizeof(float), cudaMemcpyDeviceToHost) );
|
||||||
|
|
||||||
end_t = std::chrono::steady_clock::now();
|
end_t = std::chrono::steady_clock::now();
|
||||||
std::cout << " TIME bboxes: " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
|
std::cout << " TIME bboxes: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
step_t = end_t;
|
step_t = end_t;
|
||||||
|
|
||||||
// ---------------------------------- post-process -----------------------------------------
|
// ---------------------------------- post-process -----------------------------------------
|
||||||
|
|
||||||
// --------- ctdet_post_process
|
// --------- ctdet_post_process
|
||||||
// --------- transform_preds
|
// --------- transform_preds
|
||||||
src.at<float>(0,0)=c[0];
|
|
||||||
src.at<float>(0,1)=c[1];
|
|
||||||
src.at<float>(1,0)=c[0];
|
|
||||||
src.at<float>(1,1)=c[1] + s[0] * -0.5;
|
|
||||||
dst.at<float>(0,0)=width * 0.5;
|
|
||||||
dst.at<float>(0,1)=width * 0.5;
|
|
||||||
dst.at<float>(1,0)=width * 0.5;
|
|
||||||
dst.at<float>(1,1)=width * 0.5 + width * -0.5;
|
|
||||||
|
|
||||||
src.at<float>(2,0)=src.at<float>(1,0) + (-src.at<float>(0,1)+src.at<float>(1,1) );
|
|
||||||
src.at<float>(2,1)=src.at<float>(1,1) + (src.at<float>(0,0)-src.at<float>(1,0) );
|
|
||||||
dst.at<float>(2,0)=dst.at<float>(1,0) + (-dst.at<float>(0,1)+dst.at<float>(1,1) );
|
|
||||||
dst.at<float>(2,1)=dst.at<float>(1,1) + (dst.at<float>(0,0)-dst.at<float>(1,0) );
|
|
||||||
|
|
||||||
|
|
||||||
cv::Mat trans2(cv::Size(3,2), CV_32F);
|
|
||||||
trans2 = cv::getAffineTransform( dst, src );
|
|
||||||
|
|
||||||
end_t = std::chrono::steady_clock::now();
|
|
||||||
std::cout << " TIME getAffineTrans 2: " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
|
|
||||||
step_t = end_t;
|
|
||||||
|
|
||||||
cv::Mat new_pt1(cv::Size(1,2), CV_32F);
|
cv::Mat new_pt1(cv::Size(1,2), CV_32F);
|
||||||
cv::Mat new_pt2(cv::Size(1,2), CV_32F);
|
cv::Mat new_pt2(cv::Size(1,2), CV_32F);
|
||||||
|
|
||||||
@@ -420,7 +377,7 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
|
|||||||
target_coords[i*4+2] = new_pt2.at<float>(0,0);
|
target_coords[i*4+2] = new_pt2.at<float>(0,0);
|
||||||
target_coords[i*4+3] = new_pt2.at<float>(0,1);
|
target_coords[i*4+3] = new_pt2.at<float>(0,1);
|
||||||
}
|
}
|
||||||
|
|
||||||
detected.clear();
|
detected.clear();
|
||||||
for(int i = 0; i<classes; i++){
|
for(int i = 0; i<classes; i++){
|
||||||
for(int j=0; j<K; j++)
|
for(int j=0; j<K; j++)
|
||||||
@@ -449,7 +406,7 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
|
|||||||
}
|
}
|
||||||
|
|
||||||
end_t = std::chrono::steady_clock::now();
|
end_t = std::chrono::steady_clock::now();
|
||||||
std::cout << " TIME detections: " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
|
std::cout << " TIME detections: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
|
||||||
step_t = end_t;
|
step_t = end_t;
|
||||||
|
|
||||||
std::cout<<"TOTAL: \n";
|
std::cout<<"TOTAL: \n";
|
||||||
|
|||||||
+17
-18
@@ -16,10 +16,7 @@ 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)
|
||||||
{
|
{
|
||||||
checkCuda( cudaMemcpy(topk_scores, (float *)src_begin, K*sizeof(float), cudaMemcpyDeviceToDevice) );
|
checkCuda( cudaMemcpy(topk_scores, (float *)src_begin, K*sizeof(float), cudaMemcpyDeviceToDevice) );
|
||||||
checkCuda( cudaMemcpy(topk_inds, idsrc, K*sizeof(int), cudaMemcpyDeviceToDevice) );
|
checkCuda( cudaMemcpy(topk_inds, idsrc, K*sizeof(int), cudaMemcpyDeviceToDevice) );
|
||||||
// topk_ys_[i*K +count] = (int)(ids2[j] / width);
|
|
||||||
// topk_xs_[i*K +count] = (int)(ids2[j] % width);
|
|
||||||
|
|
||||||
}
|
}
|
||||||
|
|
||||||
__global__
|
__global__
|
||||||
@@ -28,7 +25,6 @@ void sortAndTopK_kernel(dnnType *src_begin, int *idsrc, float *topk_scores, int
|
|||||||
|
|
||||||
thrust::sort_by_key(thrust::device, src_begin + i * size, src_begin + i * size + size, idsrc + i * size, thrust::greater<float>());
|
thrust::sort_by_key(thrust::device, src_begin + i * size, src_begin + i * size + size, idsrc + i * size, thrust::greater<float>());
|
||||||
thrust::copy_n(thrust::device, src_begin + i * size, K, topk_scores + i * K);
|
thrust::copy_n(thrust::device, src_begin + i * size, K, topk_scores + i * K);
|
||||||
// thrust::copy_n(thrust::device, idsrc + i * size, K, topk_inds + i * K );
|
|
||||||
thrust::copy_n(thrust::device, idsrc + i * size, K, topk_inds + i * K );
|
thrust::copy_n(thrust::device, idsrc + i * size, K, topk_inds + i * K );
|
||||||
}
|
}
|
||||||
|
|
||||||
@@ -41,20 +37,23 @@ void sortAndTopKonDevice(dnnType *src_begin, int *idsrc, float *topk_scores, int
|
|||||||
|
|
||||||
}
|
}
|
||||||
|
|
||||||
struct threshold : public thrust::binary_function<float,float,float>
|
__global__
|
||||||
{
|
void normalize_kernel(float *bgr, const int dim, const float *mean, const float *stddev){
|
||||||
__host__ __device__
|
int i = blockDim.x*blockIdx.x + threadIdx.x;
|
||||||
float operator()(float x, float y) {
|
int j = blockIdx.y;
|
||||||
double toll = 1e-6;
|
bgr[j*(dim)+i] = bgr[j*(dim)+i] - mean[j];
|
||||||
if(fabsf(x-y)>toll)
|
bgr[j*(dim)+i] = bgr[j*(dim)+i] / stddev[j];
|
||||||
return 0.0f;
|
|
||||||
else
|
}
|
||||||
return x;
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
void subtractWithThreshold(dnnType *src_begin, dnnType *src_end, dnnType *src2_begin, dnnType *src_out){
|
void normalize(float *bgr, const int ch, const int h, const int w, const float *mean, const float *stddev)
|
||||||
struct threshold op;
|
{
|
||||||
|
int num_thread = 256;
|
||||||
|
dim3 dimBlock(h*w/num_thread, ch);
|
||||||
|
normalize_kernel<<<dimBlock, num_thread, 0>>>(bgr, h*w, mean, stddev);
|
||||||
|
}
|
||||||
|
|
||||||
|
void subtractWithThreshold(dnnType *src_begin, dnnType *src_end, dnnType *src2_begin, dnnType *src_out, struct threshold op){
|
||||||
thrust::transform(thrust::device, src_begin, src_end, src2_begin, src_out, op);
|
thrust::transform(thrust::device, src_begin, src_end, src2_begin, src_out, op);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|||||||
Reference in New Issue
Block a user