From 9007e25a0064cd5916ab35e409877dfaf4000e82 Mon Sep 17 00:00:00 2001 From: Davide Sapienza Date: Fri, 7 Feb 2020 14:50:05 +0100 Subject: [PATCH] Remove mallocs and frees from the kernels Signed-off-by: Davide Sapienza --- include/sorting.h | 6 ++++-- include/tkDNN/CenternetDetection.h | 4 +++- src/CenternetDetection.cpp | 13 +++++++++---- src/sorting.cu | 20 ++++++-------------- 4 files changed, 22 insertions(+), 21 deletions(-) diff --git a/include/sorting.h b/include/sorting.h index d153c61..9c968f3 100644 --- a/include/sorting.h +++ b/include/sorting.h @@ -15,5 +15,7 @@ void topk(dnnType *src_begin, int *idsrc, int K, float *topk_scores, 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 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, float *xs_begin, float *ys_begin, dnnType *src_begin); -void bboxes(int * ids_begin, const int K, const int size, float *xs_begin, float *ys_begin, dnnType *src_begin, float *bbx0, float *bbx1, float *bby0, float *bby1); +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); +void bboxes(int * ids_begin, const int K, const int size, float *xs_begin, float *ys_begin, + dnnType *src_begin, float *bbx0, float *bbx1, float *bby0, float *bby1, float *src_out, int *ids_out); diff --git a/include/tkDNN/CenternetDetection.h b/include/tkDNN/CenternetDetection.h index 1868839..e39644f 100644 --- a/include/tkDNN/CenternetDetection.h +++ b/include/tkDNN/CenternetDetection.h @@ -75,7 +75,9 @@ class CenternetDetection { int K = 100; int width = 128;//56; // TODO - + // pointer used in the kernels + float *src_out; + int *ids_out; public: dnnType *rt_out[4]; diff --git a/src/CenternetDetection.cpp b/src/CenternetDetection.cpp index c6c4bc7..c707c6c 100644 --- a/src/CenternetDetection.cpp +++ b/src/CenternetDetection.cpp @@ -96,8 +96,13 @@ bool CenternetDetection::init(std::string tensor_path) { mean << 0.408, 0.447, 0.47; stddev << 0.289, 0.274, 0.278; - // mean << 0.485, 0.456, 0.406; - // stddev << 0.229, 0.224, 0.225; + + // Alloc array used in the kernel + checkCuda( cudaMalloc(&src_out, K *sizeof(float)) ); + checkCuda( cudaMalloc(&ids_out, K *sizeof(int)) ); + // checkCuda( cudaFree(src_out) ); + // checkCuda( cudaFree(ids_out) ); + } void CenternetDetection::testdog() { @@ -349,14 +354,14 @@ void CenternetDetection::update(cv::Mat &imageORIG) { // ----------- topk end - topKxyAddOffset(topk_inds_d, K, dim_reg.h*dim_reg.w, inttopk_xs_d, inttopk_ys_d, topk_xs_d, topk_ys_d, rt_out[3]); + topKxyAddOffset(topk_inds_d, K, dim_reg.h*dim_reg.w, inttopk_xs_d, inttopk_ys_d, topk_xs_d, topk_ys_d, rt_out[3], src_out, ids_out); // checkCuda( cudaDeviceSynchronize() ); end_t = std::chrono::steady_clock::now(); std::cout << " TIME add offset: " << std::chrono::duration_cast(end_t - step_t).count() << " ms" << std::endl; 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); + 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); // checkCuda( cudaDeviceSynchronize() ); checkCuda( cudaMemcpy(bbx0, bbx0_d, K * sizeof(float), cudaMemcpyDeviceToHost) ); diff --git a/src/sorting.cu b/src/sorting.cu index fd62cec..3ba5ac3 100644 --- a/src/sorting.cu +++ b/src/sorting.cu @@ -66,31 +66,25 @@ void topKxyclasses(int *ids_begin, int *ids_end, const int K, const int size, co } -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; - checkCuda( cudaMalloc(&src_out, K *sizeof(float)) ); +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){ thrust::gather(thrust::device, ids_begin, ids_begin + K, src_begin, src_out); thrust::transform(thrust::device, intxs_begin, intxs_begin + K, src_out, xs_begin, thrust::plus()); - int *ids_out; - checkCuda( cudaMalloc(&ids_out, K *sizeof(int)) ); thrust::transform(thrust::device, ids_begin, ids_begin + K, thrust::make_constant_iterator(size), ids_out, thrust::plus()); thrust::gather(thrust::device, ids_out, ids_out+K, src_begin, src_out); thrust::transform(thrust::device, intys_begin, intys_begin + K, src_out, ys_begin, thrust::plus()); - checkCuda( cudaFree(src_out) ); - checkCuda( cudaFree(ids_out) ); } -void bboxes(int * ids_begin, const int K, const int size, float *xs_begin, float *ys_begin, dnnType *src_begin, float *bbx0, float *bbx1, float *bby0, float *bby1){ - float *src_out; - checkCuda( cudaMalloc(&src_out, K *sizeof(float)) ); +void bboxes(int * ids_begin, const int K, const int size, float *xs_begin, float *ys_begin, + dnnType *src_begin, float *bbx0, float *bbx1, float *bby0, float *bby1, + float *src_out, int *ids_out){ thrust::gather(thrust::device, ids_begin, ids_begin + K, src_begin, src_out); thrust::transform(thrust::device, src_out, src_out + K, thrust::make_constant_iterator(2), src_out, thrust::divides()); // x0 thrust::transform(thrust::device, xs_begin, xs_begin + K, src_out, bbx0, thrust::minus()); // x1 thrust::transform(thrust::device, xs_begin, xs_begin + K, src_out, bbx1, thrust::plus()); - int *ids_out; - checkCuda( cudaMalloc(&ids_out, K *sizeof(int)) ); thrust::transform(thrust::device, ids_begin, ids_begin + K, thrust::make_constant_iterator(size), ids_out, thrust::plus()); thrust::gather(thrust::device, ids_out, ids_out + K, src_begin, src_out); thrust::transform(thrust::device, src_out, src_out + K, thrust::make_constant_iterator(2), src_out, thrust::divides()); @@ -98,7 +92,5 @@ void bboxes(int * ids_begin, const int K, const int size, float *xs_begin, float thrust::transform(thrust::device, ys_begin, ys_begin + K, src_out, bby0, thrust::minus()); // y1 thrust::transform(thrust::device, ys_begin, ys_begin + K, src_out, bby1, thrust::plus()); - checkCuda( cudaFree(src_out) ); - checkCuda( cudaFree(ids_out) ); }