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

Remove mallocs and frees from the kernels

Browse Source 
Signed-off-by: Davide Sapienza <sapienza.dav@gmail.com>
This commit is contained in:
Davide Sapienza
2020-02-07 14:50:05 +01:00
parent 289a97d06c
commit 9007e25a00
4 changed files with 22 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/sorting.h
+4 -2
View File
@@ -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);
include/tkDNN/CenternetDetection.h
+3 -1
View File
@@ -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];
src/CenternetDetection.cpp
+9 -4
View File
@@ -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<std::chrono::milliseconds>(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) );
src/sorting.cu
+6 -14
View File
@@ -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<float>());
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<int>());
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<float>());
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<float>());
// x0
thrust::transform(thrust::device, xs_begin, xs_begin + K, src_out, bbx0, thrust::minus<float>());
// x1
thrust::transform(thrust::device, xs_begin, xs_begin + K, src_out, bbx1, thrust::plus<float>());
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<int>());
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<float>());
@@ -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<float>());
// y1
thrust::transform(thrust::device, ys_begin, ys_begin + K, src_out, bby1, thrust::plus<float>());
checkCuda( cudaFree(src_out) );
checkCuda( cudaFree(ids_out) );
}