From 61aa24c6b7716b833c8a0840dc0c47a00ca7e8ff Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Tue, 30 Jun 2020 15:19:03 +0200 Subject: [PATCH] yolov4tiny works on CUDNN Signed-off-by: Micaela Verucchi --- include/tkDNN/DarknetParser.h | 1 + include/tkDNN/Layer.h | 4 +- src/DarknetParser.cpp | 4 +- src/Route.cpp | 10 +- tests/darknet/cfg/yolo4tiny.cfg | 281 ++++++++++++++++++++++++++++++++ tests/darknet/yolo4tiny.cpp | 33 ++++ 6 files changed, 328 insertions(+), 5 deletions(-) create mode 100644 tests/darknet/cfg/yolo4tiny.cfg create mode 100644 tests/darknet/yolo4tiny.cpp diff --git a/include/tkDNN/DarknetParser.h b/include/tkDNN/DarknetParser.h index f36469b..29d1e8e 100644 --- a/include/tkDNN/DarknetParser.h +++ b/include/tkDNN/DarknetParser.h @@ -11,6 +11,7 @@ namespace tk { namespace dnn { int channels = 3; int batch_normalize=0; int groups = 1; + int group_id = 0; int filters=1; int size_x=1; int size_y=1; diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 9bd8432..bd544b2 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -509,7 +509,7 @@ public: class Route : public Layer { public: - Route(Network *net, Layer **layers, int layers_n); + Route(Network *net, Layer **layers, int layers_n, int groups = 1, int group_id = 0); virtual ~Route(); virtual layerType_t getLayerType() { return LAYER_ROUTE; }; @@ -519,6 +519,8 @@ public: static const int MAX_LAYERS = 32; Layer *layers[MAX_LAYERS]; //ids of layers to be merged int layers_n; //number of layers + int groups; + int group_id; }; diff --git a/src/DarknetParser.cpp b/src/DarknetParser.cpp index 5092595..7bbc7ba 100644 --- a/src/DarknetParser.cpp +++ b/src/DarknetParser.cpp @@ -75,6 +75,8 @@ namespace tk { namespace dnn { fields.coords = std::stoi(value); else if(name.find("groups") != std::string::npos) fields.groups = std::stoi(value); + else if(name.find("group_id") != std::string::npos) + fields.group_id = std::stoi(value); else if(name.find("scale_x_y") != std::string::npos) fields.scale_xy = std::stof(value); else if(name.find("from") != std::string::npos) @@ -148,7 +150,7 @@ namespace tk { namespace dnn { //std::cout<<"Route to "<getLayerName()<<"\n"; layers.push_back(netLayers[layerIdx]); } - netLayers.push_back(new tk::dnn::Route(net, layers.data(), layers.size())); + netLayers.push_back(new tk::dnn::Route(net, layers.data(), layers.size(), f.groups, f.group_id)); } else if(f.type == "reorg") { netLayers.push_back(new tk::dnn::Reorg(net, f.stride_x)); diff --git a/src/Route.cpp b/src/Route.cpp index 39bb14e..816566e 100644 --- a/src/Route.cpp +++ b/src/Route.cpp @@ -5,7 +5,7 @@ namespace tk { namespace dnn { -Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) { +Route::Route(Network *net, Layer **layers, int layers_n, int groups, int group_id) : Layer(net) { // copy input layers if(layers_n > MAX_LAYERS) { @@ -15,6 +15,8 @@ Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) { this->layers[i] = layers[i]; } this->layers_n = layers_n; + this->groups = groups; + this->group_id = group_id; //get dims output_dim.l = 1; @@ -32,6 +34,7 @@ Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) { output_dim.c += layers[i]->output_dim.c; } + output_dim.c /= this->groups; input_dim = output_dim; checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); @@ -49,8 +52,9 @@ dnnType* Route::infer(dataDim_t &dim, dnnType* srcData) { for(int i=0; idstData; int in_dim = layers[i]->output_dim.tot(); - checkCuda( cudaMemcpy(dstData + offset, input, in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); - offset += in_dim; + int part_in_dim = in_dim / this->groups; + checkCuda( cudaMemcpy(dstData + offset, input + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + offset += part_in_dim; } //update data dimensions diff --git a/tests/darknet/cfg/yolo4tiny.cfg b/tests/darknet/cfg/yolo4tiny.cfg new file mode 100644 index 0000000..dc6f5bf --- /dev/null +++ b/tests/darknet/cfg/yolo4tiny.cfg @@ -0,0 +1,281 @@ +[net] +# Testing +#batch=1 +#subdivisions=1 +# Training +batch=64 +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.00261 +burn_in=1000 +max_batches = 500200 +policy=steps +steps=400000,450000 +scales=.1,.1 + +[convolutional] +batch_normalize=1 +filters=32 +size=3 +stride=2 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=2 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=leaky + +[route] +layers=-1 +groups=2 +group_id=1 + +[convolutional] +batch_normalize=1 +filters=32 +size=3 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=32 +size=3 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -1,-2 + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -6,-1 + +[maxpool] +size=2 +stride=2 + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=leaky + +[route] +layers=-1 +groups=2 +group_id=1 + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -1,-2 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -6,-1 + +[maxpool] +size=2 +stride=2 + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=leaky + +[route] +layers=-1 +groups=2 +group_id=1 + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -1,-2 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -6,-1 + +[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] +size=1 +stride=1 +pad=1 +filters=255 +activation=linear + + + +[yolo] +mask = 3,4,5 +anchors = 10,14, 23,27, 37,58, 81,82, 135,169, 344,319 +classes=80 +num=6 +jitter=.3 +scale_x_y = 1.05 +cls_normalizer=1.0 +iou_normalizer=0.07 +iou_loss=ciou +ignore_thresh = .7 +truth_thresh = 1 +random=0 +resize=1.5 +nms_kind=greedynms +beta_nms=0.6 + +[route] +layers = -4 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[upsample] +stride=2 + +[route] +layers = -1, 23 + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=leaky + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=linear + +[yolo] +mask = 1,2,3 +anchors = 10,14, 23,27, 37,58, 81,82, 135,169, 344,319 +classes=80 +num=6 +jitter=.3 +scale_x_y = 1.05 +cls_normalizer=1.0 +iou_normalizer=0.07 +iou_loss=ciou +ignore_thresh = .7 +truth_thresh = 1 +random=0 +resize=1.5 +nms_kind=greedynms +beta_nms=0.6 diff --git a/tests/darknet/yolo4tiny.cpp b/tests/darknet/yolo4tiny.cpp new file mode 100644 index 0000000..d9011a8 --- /dev/null +++ b/tests/darknet/yolo4tiny.cpp @@ -0,0 +1,33 @@ +#include +#include +#include "tkdnn.h" +#include "test.h" +#include "DarknetParser.h" + +int main() { + std::string bin_path = "yolo4tiny"; + std::vector input_bins = { + bin_path + "/layers/input.bin" + }; + std::vector output_bins = { + bin_path + "/debug/layer30_out.bin", + bin_path + "/debug/layer37_out.bin" + }; + std::string wgs_path = bin_path + "/layers"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4tiny.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; + downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/iRnc4pSqmx78gJs/download"); + + // parse darknet network + tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path); + net->print(); + + //convert network to tensorRT + // tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str())); + + int ret = testInference(input_bins, output_bins, net, nullptr); + net->releaseLayers(); + delete net; + // delete netRT; + return ret; +}