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

Fix the sorting kernels used in the CenterNet pre and post-processing.

Browse Source 
This commit moves the kernels in the correct sub-directory.
It creates new header file for Thrust kernels. It splits the
kernels into two files: 'normalize.cu' contains CenterNet
pre-processing operations, 'postprocessing.cu' contains the
CenterNet post-processing operations.

Signed-off-by: Davide Sapienza <sapienza.dav@gmail.com>
This commit is contained in:
Davide Sapienza
2020-04-07 18:28:12 +02:00
parent ae876e22ee
commit cf3fbeddbd
5 changed files with 31 additions and 41 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/tkDNN/CenternetDetection.h
+1 -1
View File
@@ -11,7 +11,7 @@
#include "DetectionNN.h"
#include "sorting.h"
#include "kernelsThrust.h"
namespace tk { namespace dnn {
include/sorting.h → include/tkDNN/kernelsThrust.h
+5 -5
View File
@@ -1,5 +1,6 @@
#ifndef SORTING_H
#define SORTING_H
#ifndef KERNELSTHRUST_H
#define KERNELSTHRUST_H
#include <thrust/sort.h>
#include <thrust/execution_policy.h>
@@ -9,7 +10,6 @@
#include <thrust/gather.h>
#include <thrust/copy.h>
#include "tkdnn.h"
struct threshold : public thrust::binary_function<float,float,float>
@@ -27,7 +27,7 @@ struct threshold : public thrust::binary_function<float,float,float>
void sort(dnnType *src_begin, dnnType *src_end, int *idsrc);
void topk(dnnType *src_begin, int *idsrc, int K, float *topk_scores,
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 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);
@@ -36,4 +36,4 @@ void topKxyAddOffset(int * ids_begin, const int K, const int size, int *intxs_be
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);
#endif /*SORTING_H*/
#endif //KERNELSTHRUST_H
src/CenternetDetection.cpp
-2
View File
@@ -1,10 +1,8 @@
#include "CenternetDetection.h"
#include "CenternetDetection.h"
namespace tk { namespace dnn {
bool CenternetDetection::init(const std::string& tensor_path, const int n_classes)
{
std::cout<<(tensor_path).c_str()<<"\n";
src/kernels/normalize.cu
+16
View File
@@ -0,0 +1,16 @@
#include "kernelsThrust.h"
__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[j];
bgr[j*(dim)+i] = bgr[j*(dim)+i] / stddev[j];
}
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);
}
src/sorting.cu → src/kernels/postprocessing.cu
+9 -33
View File
@@ -1,20 +1,21 @@
#include "kernelsThrust.h"
#include "sorting.h"
void sort(dnnType *src_begin, dnnType *src_end, int *idsrc)
{
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);
}
void sort(dnnType *src_begin, dnnType *src_end, int *idsrc){
thrust::sort_by_key(thrust::device,
src_begin, src_end, idsrc,
thrust::greater<float>());
// thrust::stable_sort_by_key(thrust::device,
// src_begin, src_end, idsrc,
// thrust::greater<float>());
}
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_inds, idsrc, K*sizeof(int), cudaMemcpyDeviceToDevice) );
}
@@ -22,39 +23,15 @@ void topk(dnnType *src_begin, int *idsrc, int K, float *topk_scores,
__global__
void sortAndTopK_kernel(dnnType *src_begin, int *idsrc, float *topk_scores, int *topk_inds, float *topk_ys, float *topk_xs,const int size, const int K){
int i = blockDim.x*blockIdx.x + threadIdx.x;
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, idsrc + i * size, K, topk_inds + i * K );
}
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){
int blocks = n_classes;
int threads = 1;
sortAndTopK_kernel<<<blocks, threads, 0>>>(src_begin, idsrc, topk_scores, topk_inds, topk_ys, topk_xs, size, K);
}
__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[j];
bgr[j*(dim)+i] = bgr[j*(dim)+i] / stddev[j];
}
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);
}
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);
sortAndTopK_kernel<<<blocks, threads, 0>>>(src_begin, idsrc, topk_scores, topk_inds, topk_ys, topk_xs, size, K);
}
void topKxyclasses(int *ids_begin, int *ids_end, const int K, const int size, const int wh, int *clses, int *xs, int *ys){
@@ -62,7 +39,6 @@ void topKxyclasses(int *ids_begin, int *ids_end, const int K, const int size, co
thrust::transform(thrust::device, ids_begin, ids_end, thrust::make_constant_iterator(wh), ids_begin, thrust::modulus<int>());
thrust::transform(thrust::device, ids_begin, ids_end, thrust::make_constant_iterator(size), ys, thrust::divides<int>());
thrust::transform(thrust::device, ids_begin, ids_end, thrust::make_constant_iterator(size), xs, thrust::modulus<int>());
}
void topKxyAddOffset(int * ids_begin, const int K, const int size,