yolo alternatives

This commit is contained in:
Francesco Gatti
2017-08-30 09:12:46 +00:00
parent 747fddab3f
commit a26ef98d2d
11 changed files with 1288 additions and 34 deletions
+1
View File
@@ -7,3 +7,4 @@ build/
*.caffemodel
*.h5
*.tar.gz
*.weights
+11 -1
View File
@@ -57,12 +57,22 @@ target_link_libraries(test_mnist tkDNN)
add_executable(test_mnistRT tests/mnist/test_mnistRT.cpp)
target_link_libraries(test_mnistRT tkDNN)
## YOLO NETS
add_executable(test_yolo tests/yolo/yolo.cpp)
target_link_libraries(test_yolo tkDNN)
add_executable(test_yolo_tiny tests/yolo-tiny/yolo-tiny.cpp)
add_executable(test_yolo_tiny tests/yolo_tiny/yolo_tiny.cpp)
target_link_libraries(test_yolo_tiny tkDNN)
add_executable(test_yolo_relu tests/yolo_relu/yolo_relu.cpp)
target_link_libraries(test_yolo_relu tkDNN)
add_executable(test_yolo_224 tests/yolo_224/yolo_224.cpp)
target_link_libraries(test_yolo_224 tkDNN)
################################################################################
add_executable(test_rtinference tests/test_rtinference/rtinference.cpp)
target_link_libraries(test_rtinference tkDNN)
+47 -18
View File
@@ -1,6 +1,7 @@
#include<iostream>
#include "tkdnn.h"
#include <stdlib.h> /* srand, rand */
#include <unistd.h>
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
@@ -84,35 +85,61 @@ void compute_image( cv::Mat imageORIG,
}
int print_usage() {
std::cout<<"usage: ./detection net.rt validation_list.txt [-t <thresh>] [-s]\n"
<<" -t: set thresh value\n -s: show images as compute\n\n"
<<"> validation_list.txt format: \n"
<<" path/to/image.jpg path/to/label.txt\n"
<<"> label.txt format: \n"
<<" <object-class> <x> <y> <width> <height>\n"
<<" x and y are the box center, "
<<"all values are relative to the image size\n\n";
return 1;
}
int main(int argc, char *argv[]) {
if(argc < 3) {
std::cout<<"usage: "<<argv[0]<<" net.rt validation_list.txt\n"
<<"> validation_list.txt format: \n"
<<" path/to/image.jpg path/to/label.txt\n"
<<"> label.txt format: \n"
<<" <object-class> <x> <y> <width> <height>\n"
<<" x and y are the box center, "
<<"all values are relative to the image size\n\n";
return 1;
//params
char *tensor_path = NULL;
char *imageset_path = NULL;
float thresh = 0.3f;
bool show = false;
//parse params
int c;
while ((c = getopt (argc, argv, "t:s")) != -1) {
switch(c) {
case 't': thresh = atof(optarg); break;
case 's': show = true; break;
case '?':
return print_usage();
default: return print_usage();
}
}
if(!fileExist(argv[1]))
FatalError("unable to read serialRT file");
if(argc - optind == 2) {
tensor_path = argv[optind];
imageset_path = argv[optind+1];
} else {
std::cout<<"not enough arguments.\n";
return print_usage();
}
//end parsing
if(!fileExist(tensor_path))
FatalError("unable to read serialRT file");
//convert network to tensorRT
tkDNN::NetworkRT netRT(NULL, argv[1]);
tkDNN::RegionInterpret rI(netRT.input_dim, netRT.output_dim, 80, 4, 5, 0.3f, reg_bias);
tkDNN::NetworkRT netRT(NULL, tensor_path);
tkDNN::RegionInterpret rI(netRT.input_dim, netRT.output_dim, 80, 4, 5, thresh, reg_bias);
dnnType *input = new float[netRT.input_dim.tot()];
dnnType *output = new float[netRT.output_dim.tot()];
std::string line;
std::ifstream imageset(argv[2]);
std::ifstream imageset(imageset_path);
if(!imageset.is_open())
FatalError("could not read imageset");
float mAP = 0;
int processed_images;
@@ -193,9 +220,11 @@ int main(int argc, char *argv[]) {
<<", mAP: "<<mAP/processed_images<<"\n";
//show results
//cv::namedWindow("result");
//cv::imshow("result", img);
//cv::waitKey(1);
if(show) {
cv::namedWindow("result");
cv::imshow("result", img);
cv::waitKey(1000);
}
}
return 0;
+2 -1
View File
@@ -324,7 +324,8 @@ void RegionInterpret::showImageResult(dnnType *input_h) {
for(int i=0; i<res_boxes_n; i++) {
box bx = res_boxes[i];
cv::rectangle(color, cv::Point(bx.x, bx.y), cv::Point(bx.w, bx.h),
cv::rectangle(color, cv::Point(bx.x - bx.w/2, bx.y - bx.h/2),
cv::Point(bx.x + bx.w/2, bx.y + bx.h/2),
cv::Scalar( 0, 0, 255), 2);
}
cv::namedWindow("result");
+258
View File
@@ -0,0 +1,258 @@
[net]
# Testing
#batch=1
#subdivisions=1
# Training
batch=32
subdivisions=8
width=608
height=608
channels=3
momentum=0.9
decay=0.0005
angle=0
saturation = 1.5
exposure = 1.5
hue=.1
learning_rate=0.001
burn_in=1000
max_batches = 500200
policy=steps
steps=400000,450000
scales=.1,.1
[convolutional]
batch_normalize=1
filters=32
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=64
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=128
size=3
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=64
size=1
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=128
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=256
size=3
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=128
size=1
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=256
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=512
size=3
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=256
size=1
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=512
size=3
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=256
size=1
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=512
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=1024
size=3
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=512
size=1
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=1024
size=3
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=512
size=1
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=1024
size=3
stride=1
pad=1
activation=leaky
#######
[convolutional]
batch_normalize=1
size=3
stride=1
pad=1
filters=1024
activation=leaky
[convolutional]
batch_normalize=1
size=3
stride=1
pad=1
filters=1024
activation=leaky
[route]
layers=-9
[convolutional]
batch_normalize=1
size=1
stride=1
pad=1
filters=64
activation=leaky
[reorg]
stride=2
[route]
layers=-1,-4
[convolutional]
batch_normalize=1
size=3
stride=1
pad=1
filters=1024
activation=leaky
[convolutional]
size=1
stride=1
pad=1
filters=425
activation=linear
[region]
anchors = 0.57273, 0.677385, 1.87446, 2.06253, 3.33843, 5.47434, 7.88282, 3.52778, 9.77052, 9.16828
bias_match=1
classes=80
coords=4
num=5
softmax=1
jitter=.3
rescore=1
object_scale=5
noobject_scale=1
class_scale=1
coord_scale=1
absolute=1
thresh = .6
random=1
+258
View File
@@ -0,0 +1,258 @@
[net]
# Testing
#batch=1
#subdivisions=1
# Training
batch=64
subdivisions=16
width=224
height=224
channels=3
momentum=0.9
decay=0.0005
angle=0
saturation = 1.5
exposure = 1.5
hue=.1
learning_rate=0.001
burn_in=1000
max_batches = 500200
policy=steps
steps=400000,450000
scales=.1,.1
[convolutional]
batch_normalize=1
filters=32
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=64
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=128
size=3
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=64
size=1
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=128
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=256
size=3
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=128
size=1
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=256
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=512
size=3
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=256
size=1
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=512
size=3
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=256
size=1
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=512
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=1024
size=3
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=512
size=1
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=1024
size=3
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=512
size=1
stride=1
pad=1
activation=leaky
[convolutional]
batch_normalize=1
filters=1024
size=3
stride=1
pad=1
activation=leaky
#######
[convolutional]
batch_normalize=1
size=3
stride=1
pad=1
filters=1024
activation=leaky
[convolutional]
batch_normalize=1
size=3
stride=1
pad=1
filters=1024
activation=leaky
[route]
layers=-9
[convolutional]
batch_normalize=1
size=1
stride=1
pad=1
filters=64
activation=leaky
[reorg]
stride=2
[route]
layers=-1,-4
[convolutional]
batch_normalize=1
size=3
stride=1
pad=1
filters=1024
activation=leaky
[convolutional]
size=1
stride=1
pad=1
filters=425
activation=linear
[region]
anchors = 0.57273, 0.677385, 1.87446, 2.06253, 3.33843, 5.47434, 7.88282, 3.52778, 9.77052, 9.16828
bias_match=1
classes=80
coords=4
num=5
softmax=1
jitter=.3
rescore=1
object_scale=5
noobject_scale=1
class_scale=1
coord_scale=1
absolute=1
thresh = .6
random=1
+150
View File
@@ -0,0 +1,150 @@
#include<iostream>
#include "tkdnn.h"
const char *input_bin = "../tests/yolo_224/layers/input.bin";
const char *c0_bin = "../tests/yolo_224/layers/c0.bin";
const char *c2_bin = "../tests/yolo_224/layers/c2.bin";
const char *c4_bin = "../tests/yolo_224/layers/c4.bin";
const char *c5_bin = "../tests/yolo_224/layers/c5.bin";
const char *c6_bin = "../tests/yolo_224/layers/c6.bin";
const char *c8_bin = "../tests/yolo_224/layers/c8.bin";
const char *c9_bin = "../tests/yolo_224/layers/c9.bin";
const char *c10_bin = "../tests/yolo_224/layers/c10.bin";
const char *c12_bin = "../tests/yolo_224/layers/c12.bin";
const char *c13_bin = "../tests/yolo_224/layers/c13.bin";
const char *c14_bin = "../tests/yolo_224/layers/c14.bin";
const char *c15_bin = "../tests/yolo_224/layers/c15.bin";
const char *c16_bin = "../tests/yolo_224/layers/c16.bin";
const char *c18_bin = "../tests/yolo_224/layers/c18.bin";
const char *c19_bin = "../tests/yolo_224/layers/c19.bin";
const char *c20_bin = "../tests/yolo_224/layers/c20.bin";
const char *c21_bin = "../tests/yolo_224/layers/c21.bin";
const char *c22_bin = "../tests/yolo_224/layers/c22.bin";
const char *c23_bin = "../tests/yolo_224/layers/c23.bin";
const char *c24_bin = "../tests/yolo_224/layers/c24.bin";
const char *c26_bin = "../tests/yolo_224/layers/c26.bin";
const char *c29_bin = "../tests/yolo_224/layers/c29.bin";
const char *c30_bin = "../tests/yolo_224/layers/c30.bin";
const char *g31_bin = "../tests/yolo_224/layers/g31.bin";
const char *output_bin = "../tests/yolo_224/layers/output.bin";
int main() {
// Network layout
tkDNN::dataDim_t dim(1, 3, 224, 224, 1);
tkDNN::Network net(dim);
tkDNN::Conv2d c0 (&net, 32, 3, 3, 1, 1, 1, 1, c0_bin, true);
tkDNN::Activation a0 (&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Pooling p1 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
tkDNN::Conv2d c2 (&net, 64, 3, 3, 1, 1, 1, 1, c2_bin, true);
tkDNN::Activation a2 (&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Pooling p3 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
tkDNN::Conv2d c4 (&net, 128, 3, 3, 1, 1, 1, 1, c4_bin, true);
tkDNN::Activation a4 (&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c5 (&net, 64, 1, 1, 1, 1, 0, 0, c5_bin, true);
tkDNN::Activation a5 (&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c6 (&net, 128, 3, 3, 1, 1, 1, 1, c6_bin, true);
tkDNN::Activation a6 (&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Pooling p7 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
tkDNN::Conv2d c8 (&net, 256, 3, 3, 1, 1, 1, 1, c8_bin, true);
tkDNN::Activation a8 (&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c9 (&net, 128, 1, 1, 1, 1, 0, 0, c9_bin, true);
tkDNN::Activation a9 (&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c10(&net, 256, 3, 3, 1, 1, 1, 1, c10_bin, true);
tkDNN::Activation a10(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Pooling p11(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
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, 256, 1, 1, 1, 1, 0, 0, c13_bin, true);
tkDNN::Activation a13(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c14(&net, 512, 3, 3, 1, 1, 1, 1, c14_bin, true);
tkDNN::Activation a14(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c15(&net, 256, 1, 1, 1, 1, 0, 0, c15_bin, true);
tkDNN::Activation a15(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c16(&net, 512, 3, 3, 1, 1, 1, 1, c16_bin, true);
tkDNN::Activation a16(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Pooling p17(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
tkDNN::Conv2d c18(&net, 1024, 3, 3, 1, 1, 1, 1, c18_bin, true);
tkDNN::Activation a18(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c19(&net, 512, 1, 1, 1, 1, 0, 0, c19_bin, true);
tkDNN::Activation a19(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c20(&net, 1024, 3, 3, 1, 1, 1, 1, c20_bin, true);
tkDNN::Activation a20(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c21(&net, 512, 1, 1, 1, 1, 0, 0, c21_bin, true);
tkDNN::Activation a21(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c22(&net, 1024, 3, 3, 1, 1, 1, 1, c22_bin, true);
tkDNN::Activation a22(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c23(&net, 1024, 3, 3, 1, 1, 1, 1, c23_bin, true);
tkDNN::Activation a23(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Conv2d c24(&net, 1024, 3, 3, 1, 1, 1, 1, c24_bin, true);
tkDNN::Activation a24(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Layer *m25_layers[1] = { &a16 };
tkDNN::Route m25(&net, m25_layers, 1);
tkDNN::Conv2d c26(&net, 64, 1, 1, 1, 1, 0, 0, c26_bin, true);
tkDNN::Activation a26(&net, tkDNN::ACTIVATION_LEAKY);
tkDNN::Reorg r27(&net, 2);
tkDNN::Layer *m28_layers[2] = { &r27, &a24 };
tkDNN::Route m28(&net, m28_layers, 2);
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);
tkDNN::RegionInterpret rI(dim, g31.output_dim, 80, 4, 5, 0.6f, g31_bin);
// Load input
dnnType *data;
dnnType *input_h;
readBinaryFile(input_bin, dim.tot(), &input_h, &data);
//print network model
net.print();
//convert network to tensorRT
tkDNN::NetworkRT netRT(&net, "yolo_224.rt");
dnnType *out_data, *out_data2; // cudnn output, tensorRT output
tkDNN::dataDim_t dim1 = dim; //input dim
printCenteredTitle(" CUDNN inference ", '=', 30); {
dim1.print();
TIMER_START
out_data = net.infer(dim1, data);
TIMER_STOP
dim1.print();
}
tkDNN::dataDim_t dim2 = dim;
printCenteredTitle(" TENSORRT inference ", '=', 30); {
dim2.print();
TIMER_START
out_data2 = netRT.infer(dim2, data);
TIMER_STOP
dim2.print();
}
printCenteredTitle(" CHECK RESULTS ", '=', 30);
dnnType *out, *out_h;
int out_dim = net.getOutputDim().tot();
readBinaryFile(output_bin, out_dim, &out_h, &out);
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";
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;
}
+258
View File
@@ -0,0 +1,258 @@
[net]
# Testing
#batch=1
#subdivisions=1
# Training
batch=64
subdivisions=16
width=608
height=608
channels=3
momentum=0.9
decay=0.0005
angle=0
saturation = 1.5
exposure = 1.5
hue=.1
learning_rate=0.001
burn_in=1000
max_batches = 500200
policy=steps
steps=400000,450000
scales=.1,.1
[convolutional]
batch_normalize=1
filters=32
size=3
stride=1
pad=1
activation=relu
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=64
size=3
stride=1
pad=1
activation=relu
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=128
size=3
stride=1
pad=1
activation=relu
[convolutional]
batch_normalize=1
filters=64
size=1
stride=1
pad=1
activation=relu
[convolutional]
batch_normalize=1
filters=128
size=3
stride=1
pad=1
activation=relu
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=256
size=3
stride=1
pad=1
activation=relu
[convolutional]
batch_normalize=1
filters=128
size=1
stride=1
pad=1
activation=relu
[convolutional]
batch_normalize=1
filters=256
size=3
stride=1
pad=1
activation=relu
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=512
size=3
stride=1
pad=1
activation=relu
[convolutional]
batch_normalize=1
filters=256
size=1
stride=1
pad=1
activation=relu
[convolutional]
batch_normalize=1
filters=512
size=3
stride=1
pad=1
activation=relu
[convolutional]
batch_normalize=1
filters=256
size=1
stride=1
pad=1
activation=relu
[convolutional]
batch_normalize=1
filters=512
size=3
stride=1
pad=1
activation=relu
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=1024
size=3
stride=1
pad=1
activation=relu
[convolutional]
batch_normalize=1
filters=512
size=1
stride=1
pad=1
activation=relu
[convolutional]
batch_normalize=1
filters=1024
size=3
stride=1
pad=1
activation=relu
[convolutional]
batch_normalize=1
filters=512
size=1
stride=1
pad=1
activation=relu
[convolutional]
batch_normalize=1
filters=1024
size=3
stride=1
pad=1
activation=relu
#######
[convolutional]
batch_normalize=1
size=3
stride=1
pad=1
filters=1024
activation=relu
[convolutional]
batch_normalize=1
size=3
stride=1
pad=1
filters=1024
activation=relu
[route]
layers=-9
[convolutional]
batch_normalize=1
size=1
stride=1
pad=1
filters=64
activation=relu
[reorg]
stride=2
[route]
layers=-1,-4
[convolutional]
batch_normalize=1
size=3
stride=1
pad=1
filters=1024
activation=relu
[convolutional]
size=1
stride=1
pad=1
filters=425
activation=linear
[region]
anchors = 0.57273, 0.677385, 1.87446, 2.06253, 3.33843, 5.47434, 7.88282, 3.52778, 9.77052, 9.16828
bias_match=1
classes=80
coords=4
num=5
softmax=1
jitter=.3
rescore=1
object_scale=5
noobject_scale=1
class_scale=1
coord_scale=1
absolute=1
thresh = .6
random=1
+150
View File
@@ -0,0 +1,150 @@
#include<iostream>
#include "tkdnn.h"
const char *input_bin = "../tests/yolo_relu/layers/input.bin";
const char *c0_bin = "../tests/yolo_relu/layers/c0.bin";
const char *c2_bin = "../tests/yolo_relu/layers/c2.bin";
const char *c4_bin = "../tests/yolo_relu/layers/c4.bin";
const char *c5_bin = "../tests/yolo_relu/layers/c5.bin";
const char *c6_bin = "../tests/yolo_relu/layers/c6.bin";
const char *c8_bin = "../tests/yolo_relu/layers/c8.bin";
const char *c9_bin = "../tests/yolo_relu/layers/c9.bin";
const char *c10_bin = "../tests/yolo_relu/layers/c10.bin";
const char *c12_bin = "../tests/yolo_relu/layers/c12.bin";
const char *c13_bin = "../tests/yolo_relu/layers/c13.bin";
const char *c14_bin = "../tests/yolo_relu/layers/c14.bin";
const char *c15_bin = "../tests/yolo_relu/layers/c15.bin";
const char *c16_bin = "../tests/yolo_relu/layers/c16.bin";
const char *c18_bin = "../tests/yolo_relu/layers/c18.bin";
const char *c19_bin = "../tests/yolo_relu/layers/c19.bin";
const char *c20_bin = "../tests/yolo_relu/layers/c20.bin";
const char *c21_bin = "../tests/yolo_relu/layers/c21.bin";
const char *c22_bin = "../tests/yolo_relu/layers/c22.bin";
const char *c23_bin = "../tests/yolo_relu/layers/c23.bin";
const char *c24_bin = "../tests/yolo_relu/layers/c24.bin";
const char *c26_bin = "../tests/yolo_relu/layers/c26.bin";
const char *c29_bin = "../tests/yolo_relu/layers/c29.bin";
const char *c30_bin = "../tests/yolo_relu/layers/c30.bin";
const char *g31_bin = "../tests/yolo_relu/layers/g31.bin";
const char *output_bin = "../tests/yolo_relu/layers/output.bin";
int main() {
// Network layout
tkDNN::dataDim_t dim(1, 3, 608, 608, 1);
tkDNN::Network net(dim);
tkDNN::Conv2d c0 (&net, 32, 3, 3, 1, 1, 1, 1, c0_bin, true);
tkDNN::Activation a0 (&net, CUDNN_ACTIVATION_RELU);
tkDNN::Pooling p1 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
tkDNN::Conv2d c2 (&net, 64, 3, 3, 1, 1, 1, 1, c2_bin, true);
tkDNN::Activation a2 (&net, CUDNN_ACTIVATION_RELU);
tkDNN::Pooling p3 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
tkDNN::Conv2d c4 (&net, 128, 3, 3, 1, 1, 1, 1, c4_bin, true);
tkDNN::Activation a4 (&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c5 (&net, 64, 1, 1, 1, 1, 0, 0, c5_bin, true);
tkDNN::Activation a5 (&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c6 (&net, 128, 3, 3, 1, 1, 1, 1, c6_bin, true);
tkDNN::Activation a6 (&net, CUDNN_ACTIVATION_RELU);
tkDNN::Pooling p7 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
tkDNN::Conv2d c8 (&net, 256, 3, 3, 1, 1, 1, 1, c8_bin, true);
tkDNN::Activation a8 (&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c9 (&net, 128, 1, 1, 1, 1, 0, 0, c9_bin, true);
tkDNN::Activation a9 (&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c10(&net, 256, 3, 3, 1, 1, 1, 1, c10_bin, true);
tkDNN::Activation a10(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Pooling p11(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
tkDNN::Conv2d c12(&net, 512, 3, 3, 1, 1, 1, 1, c12_bin, true);
tkDNN::Activation a12(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c13(&net, 256, 1, 1, 1, 1, 0, 0, c13_bin, true);
tkDNN::Activation a13(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c14(&net, 512, 3, 3, 1, 1, 1, 1, c14_bin, true);
tkDNN::Activation a14(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c15(&net, 256, 1, 1, 1, 1, 0, 0, c15_bin, true);
tkDNN::Activation a15(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c16(&net, 512, 3, 3, 1, 1, 1, 1, c16_bin, true);
tkDNN::Activation a16(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Pooling p17(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
tkDNN::Conv2d c18(&net, 1024, 3, 3, 1, 1, 1, 1, c18_bin, true);
tkDNN::Activation a18(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c19(&net, 512, 1, 1, 1, 1, 0, 0, c19_bin, true);
tkDNN::Activation a19(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c20(&net, 1024, 3, 3, 1, 1, 1, 1, c20_bin, true);
tkDNN::Activation a20(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c21(&net, 512, 1, 1, 1, 1, 0, 0, c21_bin, true);
tkDNN::Activation a21(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c22(&net, 1024, 3, 3, 1, 1, 1, 1, c22_bin, true);
tkDNN::Activation a22(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c23(&net, 1024, 3, 3, 1, 1, 1, 1, c23_bin, true);
tkDNN::Activation a23(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c24(&net, 1024, 3, 3, 1, 1, 1, 1, c24_bin, true);
tkDNN::Activation a24(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Layer *m25_layers[1] = { &a16 };
tkDNN::Route m25(&net, m25_layers, 1);
tkDNN::Conv2d c26(&net, 64, 1, 1, 1, 1, 0, 0, c26_bin, true);
tkDNN::Activation a26(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Reorg r27(&net, 2);
tkDNN::Layer *m28_layers[2] = { &r27, &a24 };
tkDNN::Route m28(&net, m28_layers, 2);
tkDNN::Conv2d c29(&net, 1024, 3, 3, 1, 1, 1, 1, c29_bin, true);
tkDNN::Activation a29(&net, CUDNN_ACTIVATION_RELU);
tkDNN::Conv2d c30(&net, 425, 1, 1, 1, 1, 0, 0, c30_bin, false);
tkDNN::Region g31(&net, 80, 4, 5);
tkDNN::RegionInterpret rI(dim, g31.output_dim, 80, 4, 5, 0.3f, g31_bin);
// Load input
dnnType *data;
dnnType *input_h;
readBinaryFile(input_bin, dim.tot(), &input_h, &data);
//print network model
net.print();
//convert network to tensorRT
tkDNN::NetworkRT netRT(&net, "yolo_relu.rt");
dnnType *out_data, *out_data2; // cudnn output, tensorRT output
tkDNN::dataDim_t dim1 = dim; //input dim
printCenteredTitle(" CUDNN inference ", '=', 30); {
dim1.print();
TIMER_START
out_data = net.infer(dim1, data);
TIMER_STOP
dim1.print();
}
tkDNN::dataDim_t dim2 = dim;
printCenteredTitle(" TENSORRT inference ", '=', 30); {
dim2.print();
TIMER_START
out_data2 = netRT.infer(dim2, data);
TIMER_STOP
dim2.print();
}
printCenteredTitle(" CHECK RESULTS ", '=', 30);
dnnType *out, *out_h;
int out_dim = net.getOutputDim().tot();
readBinaryFile(output_bin, out_dim, &out_h, &out);
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";
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, 608, 608);
rI.showImageResult(input_h);
return 0;
}
+139
View File
@@ -0,0 +1,139 @@
[net]
Training
batch=64
subdivisions=8
# Testing
# batch=1
# subdivisions=1
width=416
height=416
channels=3
momentum=0.9
decay=0.0005
angle=0
saturation = 1.5
exposure = 1.5
hue=.1
learning_rate=0.001
burn_in=1000
max_batches = 500200
policy=steps
steps=400000,450000
scales=.1,.1
[convolutional]
batch_normalize=1
filters=16
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=32
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=64
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=128
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=256
size=3
stride=1
pad=1
activation=leaky
[maxpool]
size=2
stride=2
[convolutional]
batch_normalize=1
filters=512
size=3
stride=1
pad=1
activation=leaky
#[maxpool]
#size=2
#stride=1
[convolutional]
batch_normalize=1
filters=1024
size=3
stride=1
pad=1
activation=leaky
###########
[convolutional]
batch_normalize=1
size=3
stride=1
pad=1
filters=512
activation=leaky
[convolutional]
size=1
stride=1
pad=1
filters=425
activation=linear
[region]
anchors = 0.57273, 0.677385, 1.87446, 2.06253, 3.33843, 5.47434, 7.88282, 3.52778, 9.77052, 9.16828
bias_match=1
classes=80
coords=4
num=5
softmax=1
jitter=.2
rescore=0
object_scale=5
noobject_scale=1
class_scale=1
coord_scale=1
absolute=1
thresh = .6
random=1
@@ -1,19 +1,19 @@
#include<iostream>
#include "tkdnn.h"
const char *input_bin = "../tests/yolo-tiny/layers/input.bin";
const char *c0_bin = "../tests/yolo-tiny/layers/c0.bin";
const char *c2_bin = "../tests/yolo-tiny/layers/c2.bin";
const char *c4_bin = "../tests/yolo-tiny/layers/c4.bin";
const char *c5_bin = "../tests/yolo-tiny/layers/c5.bin";
const char *c6_bin = "../tests/yolo-tiny/layers/c6.bin";
const char *c8_bin = "../tests/yolo-tiny/layers/c8.bin";
const char *c10_bin = "../tests/yolo-tiny/layers/c10.bin";
const char *c11_bin = "../tests/yolo-tiny/layers/c11.bin";
const char *c12_bin = "../tests/yolo-tiny/layers/c12.bin";
const char *c13_bin = "../tests/yolo-tiny/layers/c13.bin";
const char *g14_bin = "../tests/yolo-tiny/layers/g14.bin";
const char *output_bin = "../tests/yolo-tiny/layers/output.bin";
const char *input_bin = "../tests/yolo_tiny/layers/input.bin";
const char *c0_bin = "../tests/yolo_tiny/layers/c0.bin";
const char *c2_bin = "../tests/yolo_tiny/layers/c2.bin";
const char *c4_bin = "../tests/yolo_tiny/layers/c4.bin";
const char *c5_bin = "../tests/yolo_tiny/layers/c5.bin";
const char *c6_bin = "../tests/yolo_tiny/layers/c6.bin";
const char *c8_bin = "../tests/yolo_tiny/layers/c8.bin";
const char *c10_bin = "../tests/yolo_tiny/layers/c10.bin";
const char *c11_bin = "../tests/yolo_tiny/layers/c11.bin";
const char *c12_bin = "../tests/yolo_tiny/layers/c12.bin";
const char *c13_bin = "../tests/yolo_tiny/layers/c13.bin";
const char *g14_bin = "../tests/yolo_tiny/layers/g14.bin";
const char *output_bin = "../tests/yolo_tiny/layers/output.bin";
int main() {
@@ -60,7 +60,7 @@ int main() {
net.print();
//convert network to tensorRT
tkDNN::NetworkRT netRT(&net, "yolo-tiny.rt");
tkDNN::NetworkRT netRT(&net, "yolo_tiny.rt");
dnnType *out_data, *out_data2; // cudnn output, tensorRT output