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

interpret

Browse Source
This commit is contained in:
Francesco Gatti
2017-08-22 01:32:59 -07:00
parent 5a52de17eb
commit 5595b8037b
8 changed files with 106 additions and 60 deletions
Download Patch File Download Diff File Expand all files Collapse all files
CMakeLists.txt
+6 -1
View File
@@ -61,9 +61,14 @@ add_executable(test_yolo_tiny tests/yolo-tiny/yolo-tiny.cpp)
target_link_libraries(test_yolo_tiny tkDNN)
add_executable(test_rtinference tests/test_rtinference/rtinference.cpp)
target_link_libraries(test_rtinference tkDNN_static)
target_link_libraries(test_rtinference tkDNN)
#install
if (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT)
set (CMAKE_INSTALL_PREFIX "${CMAKE_BINARY_DIR}/install"
CACHE PATH "default install path" FORCE)
endif()
message("install dir:" ${CMAKE_INSTALL_PREFIX})
install(DIRECTORY include/ DESTINATION include/${CMAKE_PROJECT_NAME}
FILES_MATCHING PATTERN "*.h")
install(TARGETS tkDNN DESTINATION lib)
include/Layer.h
+24 -6
View File
@@ -276,6 +276,11 @@ public:
struct box {
float x, y, w, h;
};
struct sortable_bbox {
int index;
int cl;
float **probs;
};
/**
Region layer
@@ -284,29 +289,42 @@ struct box {
class Region : public Layer {
public:
Region(Network *net, int classes, int coords, int num, float thresh, const char* fname_weights);
Region(Network *net, int classes, int coords, int num);
virtual ~Region();
virtual layerType_t getLayerType() { return LAYER_REGION; };
int classes, coords, num;
virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
};
dnnType *bias_h, *bias_d;
class RegionInterpret {
public:
RegionInterpret(dataDim_t input_dim, dataDim_t output_dim,
int classes, int coords, int num, float thresh, const char* fname_weights);
~RegionInterpret();
dataDim_t input_dim, output_dim;
dnnType *bias_h, *bias_d; //anchors
int classes, coords, num;
float thresh;
box *boxes;
float **probs;
sortable_bbox *s;
box res_boxes[256];
int res_boxes_n;
int entry_index(int batch, int location, int entry);
box get_region_box(float *x, float *biases, int n, int index, int i, int j, int w, int h, int stride);
void get_region_boxes( float *input, int w, int h, int netw, int neth, float thresh,
float **probs, box *boxes, int only_objectness,
int *map, float tree_thresh, int relative);
void correct_region_boxes(box *boxes, int n, int w, int h, int netw, int neth, int relative);
void interpretData();
void interpretData(dnnType *data_h);
void showImageResult(dnnType *input_h);
};
}
#endif //LAYER_H
src/NetworkRT.cpp
+5 -5
View File
@@ -45,7 +45,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
builderRT->setMaxBatchSize(1);
builderRT->setMaxWorkspaceSize(1 << 30);
/*
//change datatype based on system specs
if(builderRT->platformHasFastInt8()) {
BatchStream bstream({32,dim.c, dim.h, dim.w}, 32, 1);
@@ -57,7 +57,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
dtRT = DataType::kHALF;
builderRT->setHalf2Mode(true);
}
*/
//add input layer
ITensor *input = networkRT->addInput("data", dtRT,
DimsCHW{ dim.c, dim.h, dim.w});
@@ -298,7 +298,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Region *l) {
//std::cout<<"convert Region\n";
//std::cout<<"New plugin REGION\n";
IPlugin *plugin = new RegionRT(l->classes, l->coords, l->num, l->thresh);
IPlugin *plugin = new RegionRT(l->classes, l->coords, l->num);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
checkNULL(lRT);
return lRT;
@@ -338,8 +338,8 @@ public:
if(name.find("Region") == 0) {
RegionRT *r = new RegionRT(readBUF<int>(buf), //classes
readBUF<int>(buf), //coords
readBUF<int>(buf), //num
readBUF<float>(buf)); //thesh
readBUF<int>(buf)); //num
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
src/Region.cpp
+58 -34
View File
@@ -10,15 +10,13 @@
namespace tkDNN {
Region::Region(Network *net, int classes, int coords, int num, float thresh, const char* fname_weights) :
Region::Region(Network *net, int classes, int coords, int num) :
Layer(net) {
this->classes = classes;
this->coords = coords;
this->num = num;
this->thresh = thresh;
this->res_boxes_n = 0;
// same
output_dim.n = input_dim.n;
output_dim.c = input_dim.c;
@@ -26,9 +24,6 @@ Region::Region(Network *net, int classes, int coords, int num, float thresh, con
output_dim.w = input_dim.w;
output_dim.l = input_dim.l;
//load anchors
readBinaryFile(fname_weights, 2*num, &bias_h, &bias_d);
checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) );
}
@@ -36,10 +31,12 @@ Region::~Region() {
checkCuda( cudaFree(dstData) );
}
int Region::entry_index(int batch, int location, int entry) {
int entry_index(int batch, int location, int entry,
int coords, int classes, dataDim_t &input_dim, dataDim_t &output_dim) {
int n = location / (input_dim.w*input_dim.h);
int loc = location % (input_dim.w*input_dim.h);
return batch*output_dim.tot() + n*input_dim.w*input_dim.h*(coords+classes+1) + entry*input_dim.w*input_dim.h + loc;
return batch*output_dim.tot() + n*input_dim.w*input_dim.h*(coords+classes+1) +
entry*input_dim.w*input_dim.h + loc;
}
dnnType* Region::infer(dataDim_t &dim, dnnType* srcData) {
@@ -48,16 +45,16 @@ dnnType* Region::infer(dataDim_t &dim, dnnType* srcData) {
for (int b = 0; b < dim.n; ++b){
for(int n = 0; n < num; ++n){
int index = entry_index(b, n*dim.w*dim.h, 0);
int index = entry_index(b, n*dim.w*dim.h, 0, coords, classes, input_dim, output_dim);
activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h);
index = entry_index(b, n*dim.w*dim.h, coords);
index = entry_index(b, n*dim.w*dim.h, coords, coords, classes, input_dim, output_dim);
activationLOGISTICForward(srcData + index, dstData + index, dim.w*dim.h);
}
}
//softmax start
int index = entry_index(0, 0, coords + 1);
int index = entry_index(0, 0, coords + 1, coords, classes, input_dim, output_dim);
softmaxForward(srcData + index, classes, output_dim.n*num, output_dim.tot()/num,
output_dim.w*output_dim.h, 1, output_dim.w*output_dim.h, 1, dstData + index);
@@ -66,7 +63,42 @@ dnnType* Region::infer(dataDim_t &dim, dnnType* srcData) {
}
box Region::get_region_box(float *x, float *biases, int n, int index, int i, int j, int w, int h, int stride)
/* Intepret class */
RegionInterpret::RegionInterpret(dataDim_t input_dim, dataDim_t output_dim,
int classes, int coords, int num, float thresh, const char* fname_weights) {
this->input_dim = input_dim;
this->output_dim = output_dim;
this->classes = classes;
this->coords = coords;
this->num = num;
this->thresh = thresh;
this->res_boxes_n = 0;
int tot = output_dim.w*output_dim.h*num;
boxes = (box*) malloc(tot*sizeof(box));
probs = (float**) malloc(tot*sizeof(float *));
for(int j = 0; j < tot; ++j) probs[j] = (float*) malloc((classes + 1)*sizeof(float *));
s = (sortable_bbox*) malloc(tot*sizeof(sortable_bbox));
//load anchors
readBinaryFile(fname_weights, 2*num, &bias_h, &bias_d);
}
RegionInterpret::~RegionInterpret() {
delete [] boxes;
for(int j = 0; j < output_dim.w*output_dim.h*num; ++j)
delete [] probs[j];
delete [] probs;
delete [] s;
delete [] bias_h;
checkCuda( cudaFree(bias_d) );
}
box RegionInterpret::get_region_box(float *x, float *biases, int n, int index, int i, int j, int w, int h, int stride)
{
box b;
b.x = (i + x[index + 0*stride]) / w;
@@ -76,7 +108,7 @@ box Region::get_region_box(float *x, float *biases, int n, int index, int i, int
return b;
}
void Region::get_region_boxes( float *input, int w, int h, int netw, int neth, float thresh,
void RegionInterpret::get_region_boxes( float *input, int w, int h, int netw, int neth, float thresh,
float **probs, box *boxes, int only_objectness,
int *map, float tree_thresh, int relative) {
int lh = output_dim.h;
@@ -92,14 +124,17 @@ void Region::get_region_boxes( float *input, int w, int h, int netw, int neth,
for(int j = 0; j < classes; ++j){
probs[index][j] = 0;
}
int obj_index = entry_index(0, n*lw*lh + i, coords);
int box_index = entry_index(0, n*lw*lh + i, 0);
int obj_index = entry_index(0, n*lw*lh + i,
coords, coords, classes, output_dim, output_dim);
int box_index = entry_index(0, n*lw*lh + i, 0,
coords, classes, output_dim, output_dim);
float scale = predictions[obj_index];
boxes[index] = get_region_box(predictions, bias_h, n, box_index, col, row, lw, lh, lw*lh);
float max = 0;
for(int j = 0; j < classes; ++j){
int class_index = entry_index(0, n*lw*lh + i, coords + 1 + j);
int class_index = entry_index(0, n*lw*lh + i, coords + 1 + j,
coords, classes, output_dim, output_dim);
float prob = scale*predictions[class_index];
probs[index][j] = (prob > thresh) ? prob : 0;
if(prob > max) max = prob;
@@ -111,7 +146,7 @@ void Region::get_region_boxes( float *input, int w, int h, int netw, int neth,
}
void Region::correct_region_boxes(box *boxes, int n, int w, int h, int netw, int neth, int relative) {
void RegionInterpret::correct_region_boxes(box *boxes, int n, int w, int h, int netw, int neth, int relative) {
int i;
int new_w=0;
int new_h=0;
@@ -142,11 +177,6 @@ void Region::correct_region_boxes(box *boxes, int n, int w, int h, int netw, int
//############################ BOX PROBABILITY UTILS ############################
struct sortable_bbox {
int index;
int cl;
float **probs;
};
int nms_comparator(const void *pa, const void *pb) {
sortable_bbox a = *(sortable_bbox *)pa;
sortable_bbox b = *(sortable_bbox *)pb;
@@ -196,21 +226,15 @@ int max_index(float *a, int n) {
void Region::interpretData() {
void RegionInterpret::interpretData(dnnType *data_h) {
int imW = net->input_dim.w, imH = net->input_dim.h;
int imW = input_dim.w, imH = input_dim.h;
int tot = output_dim.w*output_dim.h*num;
float *lel = new dnnType[output_dim.tot()];
cudaMemcpy(lel, dstData, output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost);
box *boxes = (box*) calloc(tot, sizeof(box));
float **probs = (float**) calloc(tot, sizeof(float *));
for(int j = 0; j < tot; ++j) probs[j] = (float*)calloc(classes + 1, sizeof(float *));
get_region_boxes(lel, imW, imH, output_dim.w, output_dim.h, thresh, probs, boxes, 0, 0, 0.5, 1);
get_region_boxes(data_h, imW, imH, output_dim.w, output_dim.h, thresh, probs, boxes, 0, 0, 0.5, 1);
//delete repeats
sortable_bbox *s = (sortable_bbox*)calloc(tot, sizeof(sortable_bbox));
for(int i = 0; i < tot; ++i){
s[i].index = i;
s[i].cl = classes;
@@ -230,13 +254,13 @@ void Region::interpretData() {
}
}
}
free(s);
//
//print results
for(int i = 0; i < tot; ++i){
int cl = max_index(probs[i], classes);
float prob = probs[i][cl];
if(prob > thresh) {
box b = boxes[i];
int x = (b.x-b.w/2.)*imW;
@@ -255,7 +279,7 @@ void Region::interpretData() {
}
}
void Region::showImageResult(dnnType *input_h) {
void RegionInterpret::showImageResult(dnnType *input_h) {
#ifdef OPENCV
dataDim_t dim = net->input_dim;
src/pluginsRT/RegionRT.cpp
+1 -4
View File
@@ -4,12 +4,11 @@
class RegionRT : public IPlugin {
public:
RegionRT(int classes, int coords, int num, float thresh) {
RegionRT(int classes, int coords, int num) {
this->classes = classes;
this->coords = coords;
this->num = num;
this->thresh = thresh;
}
~RegionRT(){
@@ -78,7 +77,6 @@ public:
tkDNN::writeBUF(buf, classes);
tkDNN::writeBUF(buf, coords);
tkDNN::writeBUF(buf, num);
tkDNN::writeBUF(buf, thresh);
tkDNN::writeBUF(buf, c);
tkDNN::writeBUF(buf, h);
tkDNN::writeBUF(buf, w);
@@ -86,7 +84,6 @@ public:
int c, h, w;
int classes, coords, num;
float thresh;
int entry_index(int batch, int location, int entry, int batchSize) {
int n = location / (w*h);
src/utils.cpp
+2 -1
View File
@@ -86,7 +86,8 @@ int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) {
int diffs = 0;
for(int i=0; i<size; i++) {
if(fabs(data_h[i] - correct_h[i]) > eps) {
if(data_h[i] != data_h[i] && correct_h[i] != correct_h[i] && //nan control
fabs(data_h[i] - correct_h[i]) > eps) {
diffs += 1;
if(diffs == 1)
std::cout<<"\n";
tests/yolo-tiny/yolo-tiny.cpp
+1 -5
View File
@@ -49,7 +49,7 @@ int main() {
tkDNN::Conv2d c12(&net, 512, 3, 3, 1, 1, 1, 1, c12_bin, true);
tkDNN::Activation a12(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c13(&net, 425, 1, 1, 1, 1, 0, 0, c13_bin, false);
tkDNN::Region g14(&net, 80, 4, 5, 0.6f, g14_bin);
tkDNN::Region g14(&net, 80, 4, 5);
// Load input
dnnType *data;
@@ -89,9 +89,5 @@ int main() {
std::cout<<"CUDNN vs correct"; checkResult(out_dim, out_data, out);
std::cout<<"TRT vs correct"; checkResult(out_dim, out_data2, out);
std::cout<<"CUDNN vs TRT "; checkResult(out_dim, out_data, out_data2);
std::cout<<"\n\nDetected objects: \n";
g14.interpretData();
g14.showImageResult(input_h);
return 0;
}
tests/yolo/yolo.cpp
+9 -4
View File
@@ -97,7 +97,9 @@ int main() {
tkDNN::Conv2d c29(&net, 1024, 3, 3, 1, 1, 1, 1, c29_bin, true);
tkDNN::Activation a29(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c30(&net, 425, 1, 1, 1, 1, 0, 0, c30_bin, false);
tkDNN::Region g31(&net, 80, 4, 5, 0.6f, g31_bin);
tkDNN::Region g31(&net, 80, 4, 5);
tkDNN::RegionInterpret rI(dim, g31.output_dim, 80, 4, 5, 0.6f, g31_bin);
// Load input
dnnType *data;
@@ -120,7 +122,7 @@ int main() {
TIMER_STOP
dim1.print();
}
tkDNN::dataDim_t dim2 = dim;
printCenteredTitle(" TENSORRT inference ", '=', 30); {
dim2.print();
@@ -139,7 +141,10 @@ int main() {
std::cout<<"CUDNN vs TRT "; checkResult(out_dim, out_data, out_data2);
std::cout<<"\n\nDetected objects: \n";
g31.interpretData();
g31.showImageResult(input_h);
dnnType *output_h = new dnnType[rI.output_dim.tot()];
checkCuda(cudaMemcpy(output_h, out_data2,
rI.output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost));
rI.interpretData(output_h);
rI.showImageResult(input_h);
return 0;
}