200 Commits

Author SHA1 Message Date
Micaela Verucchi 615b4c8a52 Udpate tracker
Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>
2020-02-27 18:04:52 +01:00
Micaela Verucchi fe37bcdf1e Merge with master
Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>
2020-02-27 17:41:26 +01:00
Micaela Verucchi 6911752e3e Update tracker
Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>
2020-02-27 17:00:07 +01:00
Francesco Gatti 443691414a LSTM ok 2020-02-16 17:21:58 +01:00
Francesco Gatti 1a1c54f364 structure ok, result wrong 2020-02-16 17:08:19 +01:00
Francesco Gatti 10b7160677 works but it need cleaning 2020-02-16 16:28:39 +01:00
Francesco Gatti 4746121d43 LSTM params 2020-02-15 20:37:08 +01:00
Francesco Gatti 4fa5d2c231 lstm return seq 2020-02-13 23:21:28 +01:00
Francesco Gatti c1c2173e4d removed unused var 2020-02-13 23:10:48 +01:00
Francesco Gatti 03d39d991c LSTM to be tested 2020-02-13 23:04:29 +01:00
Francesco Gatti a9c0db0bf6 LSTM cudnn test 2020-02-13 19:27:18 +01:00
luca 8a4d1cac17 compile with tensorrt 5 2020-01-20 14:51:38 +01:00
Francesco Gatti 146e144249 Update README.md 2020-01-16 18:24:35 +01:00
Francesco Gatti 2f57ba1222 Update README.md 2020-01-16 18:21:34 +01:00
xavier da4f246157 add DLA, plugin for shortcut and leaky. new verison 0.4 2020-01-15 21:48:18 +01:00
xavier 33844c1ab2 Batchnorm eps fix, works on jetpack 4.3 2020-01-15 19:11:14 +01:00
xavier 7233b065a8 Compiles with opencv4 -pt 2 2020-01-15 19:07:05 +01:00
xavier 2fa9f691ab Compiles with opencv4
Signed-off-by: xavier <micaelaverucchi@gmail.com>
2020-01-15 19:01:23 +01:00
Francesco Gatti f3f5daf3db Merge branch 'master' of https://github.com/ceccocats/tkDNN 2020-01-15 18:07:44 +01:00
Francesco Gatti c2d73623e5 support clion 2020-01-15 18:07:40 +01:00
xavier c32a0be257 Batchnorm eps fix, works on jetpack 4.3 2020-01-15 18:06:02 +01:00
xavier 57d7743f7e Change opencv funcion call (due to OpenCV 4)
Signed-off-by: xavier <micaelaverucchi@gmail.com>
2020-01-15 09:55:10 +01:00
mbosi 6bf9179acc fix to drivework global path 2019-12-12 12:30:24 +01:00
Francesco Gatti b218b18a02 readme update 2019-12-02 20:24:12 +01:00
Davide Sapienza cfb457fdec Fix bug in velocity conversion.
This commit fixes a wrong operation in the velocity conversion.
A reduced speed (because we are in a urban track) is now stored
into a uint8. Thus granularity is now half km/h.

Signed-off-by: Davide Sapienza <sapienza.dav@gmail.com>
2019-11-11 18:07:03 +01:00
Francesco Gatti aa5927d8a1 findCUDNN 2019-11-06 14:04:23 +01:00
Davide Sapienza 1af2b792b8 Fix segmentation fault on yolo3Detection object copy
This commit fixes a segmentation fault appeared in yolo3 network
updating. Now in Camera_t structure type there is a network pointer.

Signed-off-by: Davide Sapienza <sapienza.dav@gmail.com>
2019-10-18 19:05:42 +02:00
Davide Sapienza d5ae26dfef Add configuration file to read CLASS parameters
This commit changes the parameters reading. It introduces
getopt to read input parameters from command line and it
uses a configuration yaml file to read the input parameters
for the network, the map and the cameras.

This commit fixes a bug in message sending to the aggregator.

Signed-off-by: Davide Sapienza <sapienza.dav@gmail.com>
2019-10-15 10:24:43 +02:00
Micaela Verucchi 35787cc771 Refactoring and modularization
Signed-off-by: Micaela Verucchi <micaela.verucchi@unimore.it>
2019-10-04 11:12:01 +02:00
mbosi 92f3d1c548 fixed install cmake 2019-10-01 18:47:31 +02:00
Davide Sapienza bb7d382d96 Handle cases captured by asserts
This commit removes asserts and handles their cases.
2019-09-27 15:37:14 +02:00
Davide Sapienza 4a7b290fdf Updating the submodule 'tracker_CLASS' to the latest version 2019-09-27 14:58:26 +02:00
Davide Sapienza 5a1c7fb83d Move several cameras into a single process.
This commit pairs a camera with a thread. In this way, the
single process can to manage several camera. Every camera
thread create one video capture thread to read its input
stream. Only one camera thread can start the visualization.

This commit fixes the wrong data reading from file of the
read_projection_matrix function.

This commit fixes the wrong orientation mask accessing of
the addRoadUserfromTracker function. In the code there are
two sections to test.
2019-09-26 16:39:16 +02:00
Francesco Gatti bbc4dda635 removed buildtype 2019-09-17 17:16:46 +02:00
Francesco Gatti de8b02fe50 install fix 2019-09-17 16:12:59 +02:00
Francesco Gatti ca62784f57 include dir fix, cmake dir 2019-09-17 15:22:39 +02:00
Francesco Gatti ec02c7292f save layer names in rt file 2019-09-16 19:41:59 +02:00
Francesco Gatti 77f031c0f4 save video result 2019-09-16 10:35:29 +02:00
mbosi a038e966d9 yolo3 flir ok 2019-09-15 16:19:30 +02:00
mbosi 8c629ebe7b string input and flir test 2019-09-14 19:03:13 +02:00
Davide Sapienza b1a3620061 Fix visualization thread
This commit splits some operations into different threads.
Some threads compute the visualization preprocessing for the
live, detection, top view and disparity visualization.
Only one thread has the role to display the different views.
2019-09-09 14:40:43 +02:00
Davide Sapienza 5bbb3f3480 Include frame disparity visualization 2019-09-03 19:08:33 +02:00
Davide Sapienza 88e0f9393a Include flag to save preprocessed images 2019-09-03 19:03:55 +02:00
Davide Sapienza 2d62d2524c Fix the visualization thread
This commit moves the computation of the visualization
into the 'showImages' function (display thread). The main
thread workload and the time consuming for each frame are
reduced.
 Please enter the commit message for your changes. Lines starting
2019-08-30 17:13:29 +02:00
Davide Sapienza 7f667af48f Add some frame filters
This commit adds some box frame filters for the edge detection
(semantic segmentation) and the frame disparity operation,
both on the single frame box and on the whole image.
2019-08-30 10:07:52 +02:00
Davide Sapienza 2bcf9ab53b Update mask images 2019-08-30 10:06:12 +02:00
Davide Sapienza ed83dfd99b Edit .gitignore: it excludes generated files 2019-08-30 09:47:58 +02:00
Francesco Gatti 041968f38a cmake fix 2019-06-29 11:08:30 +02:00
Francesco Gatti f50aa4ad1a fix cmake 2019-06-28 18:51:01 +02:00
Francesco Gatti 6656c3d0e8 fix cmake 2019-06-28 17:44:32 +02:00
mive93 5f444825ad optimized undistortion 2019-05-16 19:48:59 +02:00
mive93 12fd8d0109 tracker modified 2019-05-16 12:49:23 +02:00
Micaela Verucchi 32b6d51949 update readme with dependencies 2019-05-15 09:18:40 +02:00
mive93 86da302163 added send of trackers infos 2019-05-14 08:38:52 +02:00
Autochaffeur 4ebbb6af2b README update 2019-05-13 17:32:18 +02:00
mive93 ff5e376873 added file for cameras calibration 2019-05-13 15:12:02 +02:00
mive93 a6d19d3698 order 2019-05-08 20:08:20 +02:00
mive93 e38d8e82ca new send and submodule masa_protocol added 2019-05-08 11:34:18 +02:00
mive93 caf4ddbce2 merge with master 2019-05-08 10:24:22 +02:00
mive93 8627c5feeb reading from yaml file 2019-05-07 22:43:39 +02:00
mive93 be31ae10d2 calibration 2019-05-07 22:13:24 +02:00
mive93 f51a35ac5a Merge branch 'class' of https://github.com/ceccocats/tkDNN into class 2019-05-07 21:07:23 +02:00
mive93 428858eaae mask 2019-05-07 21:07:18 +02:00
Micaela Verucchi 68ecd15125 Merge branch 'class' of https://github.com/ceccocats/tkDNN into class 2019-05-07 18:57:07 +00:00
Micaela Verucchi 3311196edb commit submodule 2019-05-07 18:56:36 +00:00
Micaela Verucchi 5c7301f7f4 Merge branch 'class' of https://github.com/ceccocats/tkDNN into class 2019-05-06 22:27:25 +02:00
Micaela Verucchi d2456b1d8a added BCDS test 2019-05-06 22:27:20 +02:00
mbosi eef1fd321f added label to demo bounding box visualization 2019-05-02 14:30:50 +02:00
Tomasz Kloda 7d1d31ac45 re-added thread for visualisation 2019-04-28 11:46:38 +00:00
Tomasz Kloda 1b9fe1ea61 added arrows(to fix), deleted old traj in top view 2019-04-27 15:59:55 +00:00
Tomasz Kloda 753699104a submodule fix 2019-04-27 11:39:41 +00:00
Tomasz Kloda 978833fd6e Revert "visualization via thread"
This reverts commit ca9d18c69e.
2019-04-27 11:27:51 +00:00
Tomasz Kloda 17c5b7a818 readme modified 2019-04-27 08:31:54 +00:00
mive93 ca9d18c69e visualization via thread 2019-04-20 17:53:03 +02:00
Micaela Verucchi 41ba8afa6d view from top added 2019-04-20 17:06:37 +02:00
Micaela Verucchi ea1f0cc193 colors to path 2019-04-20 16:19:32 +02:00
Micaela Verucchi 9b413b77ab tracking integrated 2019-04-20 15:48:52 +02:00
Micaela Verucchi 9a4a65a3c3 added data 2019-04-20 13:48:55 +02:00
Micaela Verucchi bd45b016bb minor 2019-04-20 12:37:36 +02:00
Francesco Gatti 28c012cade added tracker 2019-04-20 12:33:05 +02:00
Francesco Gatti 9b03bfcbd7 merged 2019-04-19 16:33:35 +02:00
Francesco Gatti 3c32d0c876 georeferencing 2019-04-19 16:21:56 +02:00
mive93 40c67e8536 dla commented 2019-04-15 11:54:16 +02:00
mive93 38a1956404 Merge branch 'class' of https://github.com/ceccocats/tkDNN into class 2019-04-15 11:45:33 +02:00
mive93 21a698bb63 added server and serialization 2019-04-15 11:38:54 +02:00
Francesco Gatti a85367fa22 dla commented 2019-03-07 17:40:24 +01:00
Francesco Gatti 46c32edb94 dimension inverted in tetrapack_resize test 2019-03-06 17:02:08 +01:00
Roberto Cavicchioli 3714155809 dla 2019-03-06 16:40:12 +01:00
Francesco Gatti 1c4aa3c5d7 Merge branch 'class' of https://github.com/ceccocats/tkDNN into class 2019-03-06 16:35:02 +01:00
Francesco Gatti d631169821 tetrapak test added 2019-03-06 16:29:41 +01:00
Francesco Gatti c22219ad16 DLA number print 2019-03-06 13:02:39 +01:00
mive93 61b6621d2c Updated to have more launching parameters 2019-02-22 13:02:22 +01:00
mive93 39df47574e Merge branch 'master' into class 2019-02-20 17:04:44 +01:00
mive93 1c8122f22d class stuff 2019-02-20 17:00:48 +01:00
Francesco Gatti 7505c28d2d include fix 2019-02-19 11:09:27 +00:00
rcavicchioli 39f80bbfb6 coco4 2019-02-19 11:38:13 +01:00
rcavicchioli 851c6a366c arg fix 2019-02-19 11:10:34 +01:00
Francesco Gatti de04ae1cab doc 2019-02-19 09:03:33 +00:00
Francesco Gatti 1aa4f0275d color fix 2019-02-19 08:57:51 +00:00
Francesco Gatti c7941666ec demo for more yolo3 2019-02-19 08:43:35 +00:00
Francesco Gatti 87fe342ca2 yoloRT load anchors 2019-02-18 21:39:14 +01:00
Francesco Gatti bdd8e0bc26 yolo3plug fix 2019-02-18 18:55:48 +00:00
Francesco Gatti 738fa94150 version update 2019-02-18 15:54:22 +00:00
Francesco Gatti 13063b904d yolo3 ok 2019-02-18 15:51:57 +00:00
Francesco Gatti 0d682136de yolo3 berkeley ok 2019-02-18 15:37:39 +00:00
Francesco Gatti 2c63bf05be multipl yolo morge 2019-02-06 22:24:01 +00:00
Francesco Gatti 0e97452460 dects dont works 2019-02-05 20:09:47 +00:00
Francesco Gatti c8dea4668d compute detections 2019-02-04 20:34:15 +00:00
Francesco Gatti 88097a3774 yolo3 ok 2019-01-04 22:28:10 +01:00
Francesco Gatti 2e8d0b1002 yolo3 86 route error 2018-12-23 16:20:17 +01:00
Francesco Gatti 3bd725801d upsample ok, route have problems 2018-12-22 23:56:01 +01:00
Francesco Gatti 34be4cd00f yoloRT layer 2018-12-22 21:26:50 +01:00
Francesco Gatti 3b60de00f8 2 input shortcut 2018-12-21 16:17:48 +01:00
Francesco Gatti 53b429551d 2 input shortcut 2018-12-21 16:16:38 +01:00
Francesco Gatti 64626bf547 shortcut rt test 2018-12-21 15:53:39 +01:00
Francesco Gatti 7a51b4382d yolo3 ok 2018-12-21 15:28:47 +01:00
Francesco Gatti c13bda3863 yolo layer break everything 2018-12-21 11:07:35 +01:00
Francesco Gatti 2606820300 layer 96 dont match 2018-12-20 18:17:40 +01:00
Francesco Gatti a41b22e1f2 layer 94 2018-12-20 17:35:49 +01:00
Francesco Gatti c8f2e1b448 upsample ok 2018-12-20 17:08:31 +01:00
Francesco Gatti 2ab47b5874 yolo layer 2018-12-20 16:10:01 +01:00
Francesco Gatti 67cc566a0d layer 81 2018-12-20 14:52:22 +01:00
Francesco Gatti 217ff20058 layer 61 2018-12-20 12:02:47 +01:00
Francesco Gatti 991abdb410 layer 36 2018-12-20 11:46:58 +01:00
Francesco Gatti 7a46601306 yolo3 layer 15 2018-12-20 11:36:10 +01:00
Francesco Gatti e91db28756 shortcut cu 2018-12-20 09:53:17 +01:00
Francesco Gatti ed02930464 upsample template 2018-12-19 22:45:43 +01:00
Francesco Gatti 5f25e0b5f6 shortcut template 2018-12-19 22:36:46 +01:00
Francesco Gatti bc0ea65766 yolo3 debug start 2018-12-19 19:39:31 +01:00
Francesco Gatti dc55874f14 yolo cfg 2018-12-18 18:21:56 +01:00
Francesco Gatti 70373d638b fix 2018-12-18 18:09:18 +01:00
Francesco Gatti a9970f43fb tests/yolo_berkeley/yolo_berkeley.cpp 2018-12-18 18:07:37 +01:00
Francesco Gatti 6eb63160c8 berkeley 2018-12-18 14:52:18 +01:00
Francesco Gatti 6249956469 namespace change 2018-12-14 21:55:16 +01:00
Francesco Gatti 443179359d config 2018-12-03 22:04:04 +01:00
Francesco Gatti a13bc2f007 ../CMakeLists.txt 2018-12-03 17:52:44 +01:00
Francesco Gatti 4d30f0abd7 compile on x86 2018-12-03 17:37:24 +01:00
Francesco Gatti 415bd47697 opencv include fix 2018-12-03 15:52:02 +01:00
Alessio 09679d7bb6 voc 2018-09-18 16:27:09 +02:00
Francesco Gatti 029ad71673 readme ok 2018-09-15 09:04:23 +00:00
Francesco Gatti 6331724953 live detection 2018-09-15 08:57:43 +00:00
Tomasz b7d240ea6d opencv fix 2018-09-15 08:09:00 +00:00
Francesco Gatti 2cf8d8f6fc fp16 implementation, TODO deallocate in LayerWgs 2017-08-30 14:37:25 +00:00
Francesco Gatti a26ef98d2d yolo alternatives 2017-08-30 09:12:46 +00:00
Francesco Gatti 747fddab3f usage 2017-08-29 17:04:02 +00:00
Francesco Gatti b2d6dcd207 detect demo with mAP 2017-08-29 16:48:18 +00:00
Francesco Gatti ab45c24efc check control ok 2017-08-28 00:53:39 +02:00
Francesco Gatti e449209d01 0.3 box iou thresh 2017-08-25 06:31:00 -07:00
Francesco Gatti e93ed59c30 Merge branch 'cudnn5' of https://github.com/ceccocats/tkDNN into cudnn5 2017-08-25 06:09:33 -07:00
Francesco Gatti 168a1d8b27 color 2017-08-25 06:09:29 -07:00
Francesco Gatti 6c2f6bcf2e optimization2 2017-08-25 15:07:47 +02:00
Francesco Gatti 030e14d782 spalla overlap optimization 2017-08-25 11:41:29 +02:00
Francesco Gatti 00355cfcf4 delete repeats to be optimized 2017-08-22 07:52:51 -07:00
Francesco Gatti 0119b31455 class in box 2017-08-22 06:37:43 -07:00
Francesco Gatti 37b050a9c8 opencv compile not for dw 2017-08-22 02:33:08 -07:00
Francesco Gatti c41a0a09a6 version fix 2017-08-22 01:43:36 -07:00
Francesco Gatti 5595b8037b interpret 2017-08-22 01:32:59 -07:00
Francesco Gatti 5a52de17eb driveworks compile 2017-08-21 09:50:05 -07:00
Francesco Gatti 0aa9de4ce8 better rt inference 2017-08-21 12:10:17 +00:00
Francesco Gatti c63ac6b590 install 2017-08-21 12:30:34 +02:00
Francesco Gatti 2b4b9b8e49 F16 inference 2017-08-14 10:16:29 +00:00
Francesco Gatti 66ad6bb1d6 input dim fix 2017-08-14 11:57:28 +02:00
Francesco Gatti fc9fb4f153 support check 2017-08-14 11:48:29 +02:00
Francesco Gatti 6110fffbb5 inference fix 2017-08-14 11:36:48 +02:00
Francesco Gatti b3a369dc29 RTinference test 2017-08-14 11:24:23 +02:00
Francesco Gatti 81e5f6a97b int8 2017-08-14 10:29:17 +02:00
Francesco Gatti 2d7563d27c cast fix 2017-08-11 15:20:15 +00:00
Francesco Gatti 3b2f062dd9 tensorRT serialization OK 2017-08-11 17:17:05 +02:00
Francesco Gatti 57c9a6ec99 LEAKY serialized 2017-08-11 16:32:22 +02:00
Francesco Gatti 04f96048b6 memcpyasync 2017-08-11 13:56:36 +00:00
Francesco Gatti 3124f86878 stream in TRT plugin 2017-08-10 19:21:15 +02:00
Francesco Gatti 9a6058ac4a removed sync 2017-08-10 18:47:21 +02:00
Francesco Gatti aef39f6144 opencv fix 2017-08-10 14:30:55 +00:00
Francesco Gatti 266330009c opencv viz 2017-08-10 16:22:17 +02:00
Francesco Gatti b75fa637cb better print 2017-08-09 16:04:49 +00:00
Francesco Gatti 1c6888f312 auto download 2017-08-09 14:13:17 +00:00
Francesco Gatti 3215d5aab0 tiny yolo fix 2017-08-09 12:44:06 +02:00
Francesco Gatti d7ce952465 get regions 2017-08-08 17:17:24 +02:00
Francesco Gatti 0a9957ba18 network print 2017-08-08 14:59:25 +02:00
Francesco Gatti 7d570c0df4 tiny yolo not working 2017-08-07 15:05:48 +02:00
Francesco Gatti 34198a4e8d fix 2017-08-04 16:16:03 +00:00
Francesco Gatti b20a2e2902 fix 2017-08-04 10:45:07 +02:00
Francesco Gatti 0ff47ad6ba YOLO IN TENSORT :) 2017-08-03 16:50:57 +02:00
Francesco Gatti 4e189755cf yolo weights tar 2017-08-03 16:09:23 +02:00
Francesco Gatti 858b3501fa yolo TensorRT almost DONE 2017-08-03 15:52:08 +02:00
Francesco Gatti 2ef76209a1 LEAKY plugin 2017-08-03 13:25:33 +02:00
Francesco Gatti 4526e2767a NetworkRT (deallocations to be done) 2017-08-03 12:16:57 +02:00
Francesco Gatti e8355cee67 better network model 2017-08-01 23:03:02 +02:00
Francesco Gatti 300b0af5dd mnist RT ok 2017-08-01 20:58:24 +02:00
Francesco Gatti 714bd5f757 mnist tensorrt incomplete 2017-08-01 18:58:59 +02:00
Francesco Gatti bed0b57fad mnist tensor 2017-08-01 18:08:56 +02:00
Francesco Gatti ed5e5d58b5 TensorRT version 2017-08-01 17:51:49 +02:00
Francesco Gatti 1cfe70365f yolo test 2017-08-01 17:12:29 +02:00
Francesco Gatti b94931f9f7 yolo layers 2017-08-01 16:08:56 +02:00
Francesco Gatti 8e4b3c6c17 download test data 2017-07-26 01:46:25 -09:00
Francesco Gatti 483ffefc35 better tests 2017-07-26 12:23:49 +02:00
Francesco Gatti e7a6f1fb6c Softmax 2017-07-18 16:45:19 +02:00
Francesco Gatti 0887199880 caffe wgex fix 2017-07-18 08:45:03 +02:00
Francesco Gatti e83baae1a7 caffe exporter 2017-07-17 20:18:02 +02:00
Francesco Gatti e4df86a07c cudnn5 branch for tx2 2017-07-14 11:11:51 +02:00
176 changed files with 16664 additions and 1078 deletions
+15 -1
View File
@@ -1,5 +1,19 @@
*~
demo/demo/data/img_crop/
demo/demo/data/img_disparity/
demo/demo/data/map/
demo/demo/data/masks_orient/
demo/demo/data/pmat_new/
demo/demo/data/masks_v2/
build/
.vscode/
*.bin
*.pyc
*.pyc
*.prototxt
*.caffemodel
*.h5
*.tar.gz
*.weights
*.zip
.idea/
*.hdf5
+6
View File
@@ -0,0 +1,6 @@
[submodule "tracker_CLASS"]
path = tracker_CLASS
url = https://github.com/mive93/tracker_CLASS.git
[submodule "masa_protocol"]
path = masa_protocol
url = https://git.hipert.unimore.it/rcavicchioli/masa_protocol.git
+149 -11
View File
@@ -1,17 +1,155 @@
cmake_minimum_required(VERSION 2.8)
cmake_minimum_required(VERSION 3.5)
project (tkDNN)
set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake)
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC")
include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN)
find_package(CUDA QUIET REQUIRED)
# project specific flags
if(DEBUG)
add_definitions(-DDEBUG)
endif()
cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS})
cuda_add_library(kernels SHARED src/kernels/activation_elu.cu)
include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS})
add_library(tkDNN SHARED src/Layer.cpp src/LayerWgs.cpp
src/Dense.cpp src/Activation.cpp src/Conv2d.cpp src/Conv3d.cpp src/Flatten.cpp src/MulAdd.cpp src/Pooling.cpp
src/Network.cpp src/utils.cpp)
target_link_libraries(tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDA_TOOLKIT_ROOT_DIR}/lib/libcudnn.so)
#-------------------------------------------------------------------------------
# CUDA
#-------------------------------------------------------------------------------
find_package(CUDA 9.0 REQUIRED)
SET(CUDA_SEPARABLE_COMPILATION ON)
#set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_30 --compiler-options '-fPIC'")
find_package(CUDNN REQUIRED)
# compile
file(GLOB tkdnn_CUSRC "src/kernels/*.cu")
cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIRS})
cuda_add_library(kernels SHARED ${tkdnn_CUSRC})
#-------------------------------------------------------------------------------
# External Libraries
#-------------------------------------------------------------------------------
find_package(OpenCV REQUIRED)
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV")
include_directories(/usr/include/gdal)
#-------------------------------------------------------------------------------
# Build Libraries
#-------------------------------------------------------------------------------
file(GLOB tkdnn_SRC "src/*.cpp")
set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDNN_LIBRARIES} ${OpenCV_LIBS})
file(GLOB class_SRC "src/class_src/*.cpp")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -std=c++11 -O3")
include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS} ${NVINFER_INCLUDES} "~/repos/cereal/include" ${CMAKE_CURRENT_SOURCE_DIR}/tracker_CLASS/c++/src /usr/include/python2.7)
set(class_LIBS ${OpenCV_LIBS} -lgdal yaml-cpp python2.7)
add_library(tkDNN SHARED ${tkdnn_SRC})
target_link_libraries(tkDNN ${tkdnn_LIBS})
add_library(CLASS SHARED ${class_SRC})
target_link_libraries(CLASS ${class_LIBS})
#static
#add_library(tkDNN_static STATIC ${tkdnn_SRC})
#target_link_libraries(tkDNN_static ${tkdnn_LIBS})
add_executable(test_simple tests/simple/test_simple.cpp)
target_link_libraries(test_simple tkDNN)
add_executable(test_mnist tests/mnist/test_mnist.cpp)
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_voc tests/yolo_voc/yolo_voc.cpp)
target_link_libraries(test_yolo_voc tkDNN)
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_yolo_berkeley tests/yolo_berkeley/yolo_berkeley.cpp)
target_link_libraries(test_yolo_berkeley tkDNN)
add_executable(test_yolo3_coco4 tests/yolo3_coco4/yolo3_coco4.cpp)
target_link_libraries(test_yolo3_coco4 tkDNN)
add_executable(test_yolo3_berkeley tests/yolo3_berkeley/yolo3_berkeley.cpp)
target_link_libraries(test_yolo3_berkeley tkDNN)
add_executable(test_yolo3_tetrapack tests/yolo3_tetrapack/yolo3_tetrapack.cpp)
target_link_libraries(test_yolo3_tetrapack tkDNN)
add_executable(test_yolo3_tetrapack_resize tests/yolo3_tetrapack_resize/yolo3_tetrapack_resize.cpp)
target_link_libraries(test_yolo3_tetrapack_resize tkDNN)
add_executable(test_yolo3_BCDS6 tests/yolo3_BCDS6/yolo3_BCDS6.cpp)
target_link_libraries(test_yolo3_BCDS6 tkDNN)
add_executable(test_yolo3_flir tests/yolo3_flir/yolo3_flir.cpp)
target_link_libraries(test_yolo3_flir tkDNN)
add_executable(test_imuodom tests/imuodom/imuodom.cpp)
target_link_libraries(test_imuodom tkDNN)
################################################################################
add_executable(test_rtinference tests/test_rtinference/rtinference.cpp)
target_link_libraries(test_rtinference tkDNN)
add_executable(yolo3_demo demo/demo/demo.cpp
tracker_CLASS/c++/src/ekf.cpp
tracker_CLASS/c++/src/trackutils.cpp
tracker_CLASS/c++/src/plot.cpp
tracker_CLASS/c++/src/tracker.cpp )
target_link_libraries(yolo3_demo tkDNN CLASS)
#-------------------------------------------------------------------------------
# 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/)
install(TARGETS tkDNN kernels DESTINATION lib)
install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/cmake/" # source directory
DESTINATION "share/tkDNN/cmake/" # target directory
)
#-------------------------------------------------------------------------------
# Prepare for test
#-------------------------------------------------------------------------------
set(TEST_DATA true CACHE BOOL "If true download deps")
if( ${TEST_DATA} )
message("Launching pre-build dependency installer script...")
execute_process (COMMAND bash -c "bash build_models.sh download"
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/tests)
set(TEST_DATA false CACHE BOOL "If true download deps" FORCE)
message("Finished dowloading test weights")
endif()
add_executable(tkDNNtest tests/test.cpp)
target_link_libraries(tkDNNtest tkDNN)
+33 -60
View File
@@ -1,15 +1,18 @@
# tkDNN
tkDNN is a Deep Neural Network library built with cuDNN primitives specifically thought to work on NVIDIA TK1 board.<br>
tkDNN is a Deep Neural Network library built with cuDNN primitives specifically thought to work on NVIDIA TK1(and all successive) board.<br>
The main scope is to do high performance inference on already trained models.
Currently supports the following layers:
* Dense, fully interconnected
* Activation (RELU, ELU, SIGMOID, TANH)
* Convolutional 2D
* Convolutional 3D
* Max and Average Pooling
* Flatten
* Data preprocessing
this branch actually work on every NVIDIA GPU that support the dependencies:
* CUDA 10.0
* CUDNN 7.603
* TENSORRT 6.01
* OPENCV 4.1
## Dependencies
```
sudo apt install libgdal-dev libeigen3-dev python-matplotlib libyaml-cpp-dev libcereal-dev python2.7-dev
```
## Workflow
The recommended workflow follow these step:
@@ -24,62 +27,32 @@ Build with cmake
mkdir build
cd build
cmake ..
# use -DTEST_DATA=False to skip dataset download
make
```
during the cmake configuration it will be dowloaded the weights needed for running
the tests
## Test
There is a ready to use example on *test* directory, to try it you must generate the weights with Keras
```
cd tests
python test_model.py
```
And then execute the inference on build directory
```
cd build
./tkDNNtest
```
this should output the same prediction as Keras.
Assumiung you have correctly builded the library these are the test ready to exec:
* test_simple: a simple convolutional and dense network (CUDNN only)
* test_mnist: the famous mnist netwok (CUDNN and TENSORRT)
* test_mnistRT: the mnist network hardcoded in using tensorRT apis (TENSORRT only)
* test_yolo: YOLO detection network (CUDNN and TENSORRT)
* test_yolo_tiny: smaller version of YOLO (CUDNN and TENSRRT)
* test_yolo3_berkeley: our yolo3 version trained with BDD100K dateset
## Simple example
Here is a example of the entire workflow on a simple model.
Using the following Keras model save it to a file
```python
model = Sequential()
model.add(Reshape((20, 1), input_shape=(20)))
model.add(Dense(256))
model.compile()
# save model
model.save("path/to/model.h5")
## yolo3 berkeley demo detection
For the live detection you need to precompile the tensorRT file by luncing the desidered network test, this is the recommended process:
```
After the model is created the weights can be exported for tkDNN inference
export TKDNN_MODE=FP16 # set the half floating point optimization
rm yolo3_berkeley.rt # be sure to delete(or move) old tensorRT files
./test_yolo3_berkeley # run the yolo test (is slow)
# with f16 inference the result will be a bit incorrect
```
python weights_exporter model.h5 dense --output=weights/path
this will genereate a yolo3_berkeley.rt file that can be used for live detection:
```
the exporter take as arguments, in order:
* input model
* layer type ["dense", "conv2d", conv3d"]
* { layer type ["dense", "conv2d", conv3d"] for each layer to export }
* optional argument --output define path where export weights
Then we can create a c++ program to do inference on tk1
```c++
#include<tkdnn.h> //library include
//Network object
tkDNN::Network net;
//input dimension
tkDNN::dataDim_t dim(1, 20, 1, 1, 1);
//Dense layer
tkDNN::Dense d0(&net, dim, 256, "weights/path", "bias/path");
//here load the input data to CUDA
//value_type is an alias of "float"
value_type *data_d = [...]
//do inference
value_type *output_d = d0.infer(dim, data_d);
//dim will be updated with the output dimension
```
The result is finally stored on output_d in device memory.
./yolo3_demo # launch detection on a demo video
./yolo3_demo yolo3_berkeley.rt /dev/video0 # launch detection on device 0
```
![demo](https://user-images.githubusercontent.com/11562617/72547657-540e7800-388d-11ea-83c6-49dfea2a0607.gif)
+33
View File
@@ -0,0 +1,33 @@
# Find the header files
find_path(CUDNN_INCLUDE_DIR
${CMAKE_SYSROOT}/usr/local/include
${CMAKE_SYSROOT}/usr/include
/usr/local/nvidia/tensorrt/include/
NO_DEFAULT_PATH
)
set(OLD_ROOT ${CMAKE_FIND_ROOT_PATH})
list(APPEND CMAKE_FIND_ROOT_PATH /)
list(APPEND CMAKE_FIND_LIBRARY_SUFFIXES .so.7)
list(APPEND CMAKE_FIND_LIBRARY_SUFFIXES .so.5)
find_library(CUDNN_LIB
NAMES cudnn
PATHS
/usr/local/driveworks/targets/${CMAKE_SYSTEM_PROCESSOR}-Linux/lib
/usr/lib/${CMAKE_SYSTEM_PROCESSOR}-linux-gnu/
NO_DEFAULT_PATH
)
find_library(CUDNN_NVLIB
NAMES "nvinfer"
PATHS
/usr/local/driveworks/targets/${CMAKE_SYSTEM_PROCESSOR}-Linux/lib
/usr/lib/${CMAKE_SYSTEM_PROCESSOR}-linux-gnu/
NO_DEFAULT_PATH
)
set(CMAKE_FIND_ROOT_PATH ${OLD_ROOT})
set(CUDNN_LIBRARIES ${CUDNN_LIB} ${CUDNN_NVLIB})
message("-- Found CUDNN: " ${CUDNN_LIB})
message("-- Found NVINFER: " ${CUDNN_NVLIB})
set(CUDNN_FOUND true)
+24
View File
@@ -0,0 +1,24 @@
message("-- Found tkDNN")
set(CMAKE_MODULE_PATH ${CMAKE_CURRENT_LIST_DIR})
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --std=c++11 -fPIC")
find_package(CUDA REQUIRED)
find_package(OpenCV REQUIRED)
find_package(CUDNN REQUIRED)
set(tkDNN_INCLUDE_DIRS
${CUDA_INCLUDE_DIRS}
${OPENCV_INCLUDE_DIRS}
${CUDNN_INCLUDE_DIRS}
)
set(tkDNN_LIBRARIES
tkDNN
kernels
${CUDA_LIBRARIES}
${CUDA_CUBLAS_LIBRARIES}
${CUDNN_LIBRARIES}
${OpenCV_LIBS}
)
set(tkDNN_FOUND true)
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Mon 06 May 2019 11:32:13 PM CEST"
image_width: 1920
image_height: 1080
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 1.6158690952190570e+03, 0., 9.4702812371722337e+02, 0.,
1.6123979985757153e+03, 5.1995630055718266e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -4.0971199964304100e-01, 1.8755404192050384e-01,
-5.3059322427743867e-03, -1.0380603625304912e-03, 0. ]
avg_reprojection_error: 3.4351035832972515e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Tue 07 May 2019 10:32:53 AM CEST"
image_width: 3072
image_height: 1728
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 4.7264390181579711e+03, 0., 1.5059850098280642e+03, 0.,
4.6793092340700096e+03, 6.7300681982359868e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -4.2669569210605879e-01, 6.6337608795749903e-01,
-1.3881256269106437e-03, 5.2468063845700682e-03, 0. ]
avg_reprojection_error: 3.1312290189919406e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Tue 07 May 2019 10:14:44 AM CEST"
image_width: 1920
image_height: 1080
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 1.6902498656747011e+03, 0., 9.7959318966703324e+02, 0.,
1.7552884617253583e+03, 5.3327953707582492e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -5.4891909767312119e-01, 2.5555919841568631e-01,
-4.3831358875660656e-03, -1.3934378903760349e-02, 0. ]
avg_reprojection_error: 1.1758482932800183e+00
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Mon 06 May 2019 11:32:13 PM CEST"
image_width: 1920
image_height: 1080
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 1.6158690952190570e+03, 0., 9.4702812371722337e+02, 0.,
1.6123979985757153e+03, 5.1995630055718266e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -4.0971199964304100e-01, 1.8755404192050384e-01,
-5.3059322427743867e-03, -1.0380603625304912e-03, 0. ]
avg_reprojection_error: 3.4351035832972515e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Tue 07 May 2019 09:56:50 AM CEST"
image_width: 1920
image_height: 1080
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 1.6229477302581809e+03, 0., 1.0277357980566628e+03, 0.,
1.6485741394129034e+03, 5.5596919291027621e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -3.7853584845653426e-01, 7.8553352913896368e-02,
-6.5552938633907229e-03, -1.6436824648695104e-02, 0. ]
avg_reprojection_error: 8.4629096638637347e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Fri 03 May 2019 11:56:13 PM CEST"
image_width: 3072
image_height: 1728
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 5.8796921906556563e+03, 0., 1.3036708932691290e+03, 0.,
5.9435402023228071e+03, 8.1110067822514861e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -5.4688862790206871e-01, 5.1913397860290666e-01,
-2.1076612628273591e-03, 1.6869796115416984e-02, 0. ]
avg_reprojection_error: 6.7667474319420251e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Sat 04 May 2019 12:00:38 AM CEST"
image_width: 3072
image_height: 1728
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 4.6903033136815602e+03, 0., 1.6303445000881884e+03, 0.,
4.7582671272189546e+03, 4.3596515032334111e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -3.4366857232996317e-01, 2.2799325522263861e-01,
2.0765840315530557e-02, -4.0088654509745098e-03, 0. ]
avg_reprojection_error: 3.9811872397860709e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Tue 07 May 2019 10:20:53 AM CEST"
image_width: 3072
image_height: 1728
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 2.5005410461483498e+03, 0., 1.5319405824251596e+03, 0.,
2.5001544574623872e+03, 7.8267345299919543e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -3.7379752112038928e-01, 1.6246299444310250e-01,
8.0371978716752837e-04, -9.6108499236087584e-04, 0. ]
avg_reprojection_error: 3.6334262234685299e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Tue 07 May 2019 10:37:47 AM CEST"
image_width: 960
image_height: 720
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 5.0439587680799593e+02, 0., 4.8997081391816727e+02, 0.,
5.0714582349015507e+02, 3.5481348085748095e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -2.7916140864065331e-01, 6.5465070220501562e-02,
-1.9231901334709591e-03, -2.6191562264760264e-03, 0. ]
avg_reprojection_error: 5.7283635087126605e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Tue 07 May 2019 10:41:12 AM CEST"
image_width: 960
image_height: 720
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 5.1663651913150818e+02, 0., 4.7267297458218127e+02, 0.,
5.1291090124818436e+02, 3.8505850298928243e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -2.8051872523046845e-01, 6.0895269981008610e-02,
-9.7920840355269542e-03, -4.9804820350633240e-04, 0. ]
avg_reprojection_error: 5.4967308787122626e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Tue 07 May 2019 10:50:02 AM CEST"
image_width: 960
image_height: 720
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 4.9724079419911664e+02, 0., 4.9277930193807083e+02, 0.,
4.9700744926387819e+02, 3.6581239154403062e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -2.7582961261093608e-01, 6.6908017283259263e-02,
-2.1580546593114500e-03, -1.7921711595441153e-03, 0. ]
avg_reprojection_error: 3.7129088933918375e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Tue 07 May 2019 10:53:27 AM CEST"
image_width: 960
image_height: 720
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 4.9372152821507876e+02, 0., 4.7585791077351445e+02, 0.,
4.9644139996881893e+02, 3.5961856724726260e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -2.9023109325424973e-01, 8.3150964750672046e-02,
-6.1378621304345154e-04, 8.4481910933416999e-04, 0. ]
avg_reprojection_error: 3.4691001942524069e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Sat 04 May 2019 12:35:58 AM CEST"
image_width: 3840
image_height: 2160
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 2.1723071272381276e+03, 0., 1.9718118689531000e+03, 0.,
2.2377541672328439e+03, 9.3157209524899565e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -3.8516162509048857e-01, 1.8961757063227327e-01,
1.8297248985443184e-02, -8.9166274086698288e-03, 0. ]
avg_reprojection_error: 9.9886914863900311e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Sun 05 May 2019 08:52:13 PM CEST"
image_width: 3072
image_height: 1728
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 2.9841357325808735e+03, 0., 1.5379802472694901e+03, 0.,
2.9784613885271938e+03, 8.9330228722164566e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -5.1096356919758967e-01, 1.4543132746407733e-01,
-3.1254001577433334e-02, -1.4769334036191385e-02, 0. ]
avg_reprojection_error: 9.3544537534095662e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Sun 05 May 2019 09:24:33 PM CEST"
image_width: 3072
image_height: 1728
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 2.4796388675592771e+03, 0., 1.5358283835422017e+03, 0.,
2.4440198814655632e+03, 8.9911455540136217e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -3.9478938399374452e-01, 1.6288159087710818e-01,
-1.8565610712959927e-02, -7.0112574756757643e-03, 0. ]
avg_reprojection_error: 4.5805459213906724e-01
+22
View File
@@ -0,0 +1,22 @@
%YAML:1.0
---
calibration_time: "Tue 07 May 2019 10:03:44 AM CEST"
image_width: 3072
image_height: 1728
board_width: 8
board_height: 6
square_size: 2.4799999237060547e+01
flags: 0
camera_matrix: !!opencv-matrix
rows: 3
cols: 3
dt: d
data: [ 2.6027348174982544e+03, 0., 1.4808496083807213e+03, 0.,
2.6008830910556521e+03, 6.7577068120137187e+02, 0., 0., 1. ]
distortion_coefficients: !!opencv-matrix
rows: 5
cols: 1
dt: d
data: [ -3.4912899377320661e-01, 1.5704840296202566e-01,
6.4926875404798358e-03, 5.7293259996249049e-03, 0. ]
avg_reprojection_error: 4.0040122960491076e-01
Binary file not shown.

After

Width:  |  Height:  |  Size: 5.3 MiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 3.0 MiB

Binary file not shown.
Binary file not shown.

After

Width:  |  Height:  |  Size: 28 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 91 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 86 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 36 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 40 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 38 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 37 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 45 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 39 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 82 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 97 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 89 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 86 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 103 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 101 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 15 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 18 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 16 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 17 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 16 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 23 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 16 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 17 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 128 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 142 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 82 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 86 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 82 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 61 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 86 KiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 83 KiB

@@ -0,0 +1,3 @@
-7.327951322519857 -89.11104695972499 9201.012307315486
-3.2739347567332553 -42.14373356189513 3668.5149400959867
-0.0009310035308911926 -0.010684335913469737 0.9999999999999999
@@ -0,0 +1,3 @@
-5.575118981764042 19.15613757077127 7116.209593619794
-3.2016502214672586 11.087267568165238 4660.384266971425
-0.0007553414127200226 0.002676057142671632 1.0
@@ -0,0 +1,3 @@
-0.6209787370656474 -57.192772931734694 6553.189887121747
-0.655304496315561 -28.195100762686465 3423.351597211246
-0.0001275987755308679 -0.007884524572727637 1.0
@@ -0,0 +1,3 @@
21.441694474392808 72.45500438933801 4652.4523150318
9.195810184174585 36.76655098601718 1801.6321216143722
0.0027471614789484795 0.009610371445379061 1.0
@@ -0,0 +1,3 @@
23.89713091621312 60.523880117656084 7778.591483347838
11.708054761838232 41.819183783017614 1126.9348661897782
0.0029807026524197224 0.008783281479831609 0.9999999999999999
@@ -0,0 +1,3 @@
-2.986713098145805 -21.641128011999868 6870.903103366085
-2.2937864581163416 -14.410613254705751 4594.877638262097
-0.0004752252182089383 -0.003112785793525571 1.0
@@ -0,0 +1,3 @@
0.797119441300711 35.495698816749574 2225.384878217832
0.3510443326339413 59.809425152982726 1419.4212456173686
-1.8197506232485688e-05 0.01820976131616732 1.0
@@ -0,0 +1,3 @@
2.213864892663308 -21.831117643802187 4060.61306610222
0.9075230321117119 -7.525265027425343 1268.217122301812
0.0005431549026467866 -0.0054930502425022745 1.0
@@ -0,0 +1,3 @@
-8.65061527352736 -23.247834477892475 4009.9545358742134
-2.2093309523396223 -8.398154406654037 1529.4606465341913
-0.001952580524058866 -0.006089308376531831 1.0
@@ -0,0 +1,3 @@
-8.665173353556622 -19.056833294770822 3827.027102449301
-2.728678410673174 -6.695956964089541 1523.1971944608492
-0.0020218867007343213 -0.004877696027397352 1.0
@@ -0,0 +1,3 @@
-0.5624066154780419 -26.36968028701744 3384.7920236393156
-0.6590389521668137 -8.709883019993336 1162.7049573736972
-0.00023421305979212285 -0.006824826866955653 0.9999999999999999
@@ -0,0 +1,3 @@
-0.9691844827251948 -16.63913290283819 3991.384826153086
-0.47672518258283236 -11.365071723693541 1770.7201067235435
-0.0003215600914247826 -0.005142032883105236 1.0
+3
View File
@@ -0,0 +1,3 @@
-1.5172087045737013 5.226213011053854 2493.822558690447
-1.4892268949636678 4.265359557115566 3382.185465939972
-0.000492856899218649 0.0016917382715537009 1.0
+3
View File
@@ -0,0 +1,3 @@
3.671630456405992 102.56885593938215 3570.381264275829
1.5985035927893978 46.57491080817888 -1491.4874345505793
0.0006446393888135914 0.027894339637393125 1.0
+3
View File
@@ -0,0 +1,3 @@
0.44666021794890903 -32.932455985363575 3557.76647014547
1.0444994361599629 -37.23760709667856 4404.4517853552425
0.00029114742043039743 -0.010721230871024154 1.0
+3
View File
@@ -0,0 +1,3 @@
0.1982518156851575 -13.22038303683631 595.3838470607648
-0.4806350794284377 -6.567788488139482 598.1732743332492
-0.00016595895699412456 -0.007499178419813009 1.0
+3
View File
@@ -0,0 +1,3 @@
-0.4122368700442484 -10.479982981650977 812.1821012303
-0.6536846365005352 -4.951476039607947 512.0831604767536
-0.00042630219831388434 -0.006003223898658603 1.0
Binary file not shown.
+603
View File
@@ -0,0 +1,603 @@
#include <time.h>
#include "utils.h"
#include "Yolo3Detection.h"
#include "message.h"
#include "visualization.h"
#include "configuration.h"
#include "tracker.h"
#include "../masa_protocol/include/send.hpp"
#include "../masa_protocol/include/serialize.hpp"
// #include <assert.h>
// #include <unistd.h>
// #include <mutex>
// #include <ctime>
// #include <pthread.h>
// #include <signal.h>
// #include <chrono>
// #include <math.h>
// #include <typeinfo>
// #include <iostream>
#define MAX_DETECT_SIZE 100
bool gRun;
std::chrono::steady_clock::time_point local_clock_start;
std::mutex mutexgRun;
std::string obj_class[10]{"person", "car", "truck", "bus", "motor", "bike", "rider", "traffic light", "traffic sign", "train"};
//mutex for some opencv operations
std::mutex mutex_cv;
Show_t updates;
bool SAVE_RESULT = false;
void sig_handler(int signo)
{
std::cout << "request gateway stop\n";
mutexgRun.lock();
gRun = false;
mutexgRun.unlock();
}
void *readVideoCapture(void *x_void_ptr)
{
std::cout << "readVideoCapture start...\n";
Frame_t *info_f = (Frame_t *)x_void_ptr;
mutex_cv.lock();
cv::VideoCapture cap(info_f->input, cv::CAP_FFMPEG);
mutex_cv.unlock();
cv::Mat frame_loc, frame0;
int frame_nbr_loc = 0;
// bool to_show = false;
if (!cap.isOpened())
{
mutexgRun.lock();
gRun = false;
mutexgRun.unlock();
}
else
std::cout << "camera started\n";
// cap.set(cv::CAP_PROP_BUFFERSIZE,3);
// std::cout<<"buf size: "<<cap.get(CV_CAP_PROP_BUFFERSIZE)<<std::endl;
auto start_t = std::chrono::steady_clock::now();
auto step_t = std::chrono::steady_clock::now();
auto end_t = std::chrono::steady_clock::now();
auto current_timestamp = std::chrono::steady_clock::now();
// compute fps and find camera's clock
double shift, mean_time = 0;
std::cout << "Frames per second using video.get(cv::CAP_PROP_FPS) : " << cap.get(cv::CAP_PROP_FPS) << std::endl;
std::cout << "readVideoCapture computes frame rate...\n";
// //compute frame rate
int i = 0;
int num_f = 120;
// the first 20 frames are null
while (i < 21)
{
cap >> frame_loc;
i++;
}
i = 0;
start_t = std::chrono::steady_clock::now();
while (i < num_f)
{
step_t = std::chrono::steady_clock::now();
cap >> frame_loc;
mean_time = mean_time + std::chrono::duration_cast<std::chrono::milliseconds>(std::chrono::steady_clock::now() - step_t).count();
std::cout << " step " << i << " : " << std::chrono::duration_cast<std::chrono::milliseconds>(std::chrono::steady_clock::now() - step_t).count() << " ms" << std::endl;
i++;
}
end_t = std::chrono::steady_clock::now();
std::cout << "Capturing " << num_f << " frames" << std::endl;
std::cout << " Time taken : " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - start_t).count() << " ms" << std::endl;
/*
mean_time indicates the milliseconds from a frame and the next. (frame rate)
local_clock_sync is the camera clock.
shift is the difference from local camera clock and local process clock.
a frame is allowed if its local timestamp minus its local clock is less then a tollerance,
otherwise it will be considered old.
*/
auto local_clock_sync = std::chrono::steady_clock::now();
mean_time = mean_time / num_f;
shift = ((double)std::chrono::duration_cast<std::chrono::milliseconds>(local_clock_sync - local_clock_start).count()) / mean_time;
shift = (shift - (int)shift) * mean_time;
std::cout << ".-------------------------------\n";
std::cout << " mean time: " << mean_time << std::endl;
std::cout << " shift: " << shift << std::endl;
std::cout << " TIMEDIFFERENCE: " << std::chrono::duration_cast<std::chrono::milliseconds>(local_clock_sync - local_clock_start).count() << std::endl;
std::cout << "\n\n\n\n";
std::cout << "readVideoCapture start to capture...\n";
while (gRun)
{
// mutex_cv.lock();
cap >> frame_loc;
// mutex_cv.unlock();
current_timestamp = std::chrono::steady_clock::now();
shift = std::chrono::duration_cast<std::chrono::milliseconds>(current_timestamp - local_clock_sync).count();
std::cout << " RELATIVE TIMESTAMP FRAME : " << shift << " ms" << std::endl;
shift = shift / mean_time;
shift = (shift - (int)shift) * mean_time;
shift = (shift - mean_time / 2 >= 0) ? -(mean_time - shift) : shift;
std::cout << "DELAY frame_" << frame_nbr_loc << " : " << shift << " ms" << std::endl;
// TODO: here introduce a tollerance to discard old frame
// std::cout<< "CV_CAP_PROP_POS_MSEC: "<< cap.get( cv::CAP_PROP_POS_MSEC) <<std::endl;
// std::cout<< "CV_CAP_PROP_POS_FRAMES: "<< cap.get( cv::CAP_PROP_POS_FRAMES) <<std::endl; // <-- the v4l2 'sequence' field
// std::cout<< "cv::CAP_PROP_FPS: "<< cap.get( cv::CAP_PROP_FPS)<<std::endl;
// std::cout << "Format: " << cap.get(CV_CAP_PROP_FORMAT) << "\n";
// CAP_PROP_POS_MSEC Current position of the video file in milliseconds or video capture timestamp.
std::cout << "id: " << cap.get(cv::CAP_PROP_POS_MSEC) << std::endl;
// CAP_PROP_FRAME_COUNT Number of frames in the video file.
std::cout << "id: " << cap.get(cv::CAP_PROP_FRAME_COUNT) << std::endl;
if (!frame_loc.data)
{
usleep(1000000);
mutex_cv.lock();
cap.open(info_f->input);
printf("cap reinitialize\n");
mutex_cv.unlock();
continue;
}
end_t = std::chrono::steady_clock::now();
std::cout << " VC-TIME 1 : " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - start_t).count() << " ms" << std::endl;
start_t = end_t;
info_f->sem_vc.lock();
info_f->frame = frame_loc.clone();
info_f->frame_nbr = frame_nbr_loc;
info_f->sem_vc.unlock();
// usleep(50000);
end_t = std::chrono::steady_clock::now();
std::cout << " VC-TIME 2 : " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - start_t).count() << " ms" << std::endl;
start_t = end_t;
frame_nbr_loc++;
}
return (void *)0;
}
void *computationTask(void *x_void_ptr)
{
Camera_t *camera = (Camera_t *)x_void_ptr;
pthread_t visual, originalshow, detectionshow, topviewshow, disparityshow;
pthread_t videocap;
tk::dnn::Yolo3Detection yolo = *(camera->yolo);
//create video capture thread
Frame_t info_f;
info_f.input = camera->input;
if (pthread_create(&videocap, NULL, readVideoCapture, (void *)&info_f))
{
fprintf(stderr, "Error creating thread\n");
return (void *)1;
};
bool to_show = camera->to_show;
double adfGeoTransform[6];
for (int i = 0; i < 6; i++)
adfGeoTransform[i] = camera->adfGeoTransform[i];
ModFrame_t info_show;
if (to_show)
{
// initialize updates struct
updates.update_o = false;
updates.update_de = false;
updates.update_t = false;
updates.update_di = false;
if (pthread_create(&visual, NULL, show_updates, (void *)NULL))
{
fprintf(stderr, "Error creating thread\n");
return (void *)1;
};
if (pthread_create(&originalshow, NULL, originalFrame, (void *)&info_f))
{
fprintf(stderr, "Error creating thread\n");
return (void *)1;
};
if (pthread_create(&disparityshow, NULL, disparityFrame, (void *)&info_f))
{
fprintf(stderr, "Error creating thread\n");
return (void *)1;
};
info_show.H = cv::Mat(cv::Size(3, 3), CV_64FC1);
if (pthread_create(&detectionshow, NULL, detectionFrame, (void *)&info_show))
{
fprintf(stderr, "Error creating thread\n");
return (void *)1;
};
if (pthread_create(&topviewshow, NULL, topviewFrame, (void *)&info_show))
{
fprintf(stderr, "Error creating thread\n");
return (void *)1;
};
}
char *pmatrix = camera->pmatrix;
/*projection matrix from camera to map*/
cv::Mat H(cv::Size(3, 3), CV_64FC1);
read_projection_matrix(H, pmatrix);
assert(cv::countNonZero(H) > 0);
// std::cout<<H<<std::endl;
// return (void*)0;
/*Camera calibration*/
cv::Mat cameraMat, distCoeff;
readCameraCalibrationYaml(camera->cameraCalib, cameraMat, distCoeff);
std::cout << cameraMat << std::endl;
std::cout << distCoeff << std::endl;
/*GPS information*/
std::vector<ObjCoords> coords;
/*socket*/
Communicator Comm(SOCK_DGRAM);
Comm.open_client_socket((char *)"127.0.0.1", 8888);
Message *m = new Message;
m->cam_idx = camera->CAM_IDX;
m->lights.clear();
/*Conversion for tracker, from gps to meters and viceversa*/
// mutex_cv.lock();
geodetic_converter::GeodeticConverter gc;
gc.initialiseReference(44.655540, 10.934315, 0);
// mutex_cv.unlock();
double east, north, up;
// double lat, lon, alt;
/*Mask info*/
cv::Mat mask = cv::imread(camera->maskfile, cv::IMREAD_GRAYSCALE);
cv::Mat maskOrient = cv::imread(camera->maskFileOrient);
// cv::Mat maskOrient = cv::imread(camera->maskFileOrient, 0);
/*for(int i=0; i< mask.cols; i++)
{
for(int j=0; j< mask.rows; j++)
{
std::cout<<maskOrient.at<cv::Vec3b>(i,j) <<std::endl;
}
}
return 0;*/
/*tracker infos*/
std::vector<Tracker> trackers;
std::vector<Data> cur_frame;
int initial_age = -5;
int age_threshold = -8;
int n_states = 5;
float dt = 0.03;
int frame_nbr = 0;
//save video
/*cv::VideoWriter outputVideo;
cv::Size S = cv::Size((int)cap.get(cv::CAP_PROP_FRAME_WIDTH), //Acquire input size
(int)cap.get(cv::CAP_PROP_FRAME_HEIGHT));
outputVideo.open("test.avi", static_cast<int>(cap.get(cv::CAP_PROP_FOURCC)), cap.get(cv::CAP_PROP_FPS), S, true);*/
cv::Mat map1, map2;
auto start_t = std::chrono::steady_clock::now();
auto step_t = std::chrono::steady_clock::now();
auto end_t = std::chrono::steady_clock::now();
// auto step_t_segmentation = std::chrono::steady_clock::now();
// auto end_t_segmentation = std::chrono::steady_clock::now();
//TODO: move in a thread
// // information for the disparity map
// std::vector <cv::Rect> pre_rois;
// cv::Mat pre_frame;
cv::Mat orig_frame;
// cv::Mat canny, pre_canny, canny_RGB, pre_canny_RGB;
// cv::Mat canny_img;
// box variable
tk::dnn::box b;
int x0, h, y1; //w, x1, y0;
int objClass;
std::string det_class;
;
// float prob;
cv::Scalar intensity;
// cv::VideoWriter resultVideo;
// if(SAVE_RESULT) {
// int w = cap.get(cv::CAP_PROP_FRAME_WIDTH);
// int h = cap.get(cv::CAP_PROP_FRAME_HEIGHT);
// resultVideo.open("result.mp4", cv::VideoWriter::fourcc('M','P','4','V'), 30, cv::Size(w, h));
// }
cv::Mat frame;
cv::Mat frame_crop;
cv::Mat dnn_input;
bool first_iteration = true;
while (gRun)
{
TIMER_START
start_t = std::chrono::steady_clock::now();
step_t = start_t;
info_f.sem_vc.lock();
frame = info_f.frame.clone();
if (info_f.frame_nbr - frame_nbr > 1)
std::cout << "more than one - f_n (diff " << info_f.frame_nbr - frame_nbr << ")\n";
frame_nbr = info_f.frame_nbr;
info_f.sem_vc.unlock();
std::cout << "f_n: " << frame_nbr << std::endl;
// if (!frame.data)
if (frame_nbr == 0)
{
usleep(1000000);
printf("no frame received\n");
continue;
}
orig_frame = frame.clone();
// mutex_cv.lock();
if (first_iteration)
cv::initUndistortRectifyMap(cameraMat, distCoeff, cv::Mat(), cameraMat, frame.size(), CV_16SC2, map1, map2);
cv::Mat temp = frame.clone();
cv::remap(temp, frame, map1, map2, 1);
//undistort(temp, frame, cameraMat, distCoeff);
// mutex_cv.unlock();
// this will be resized to the net format
dnn_input = frame.clone();
// TODO: async infer
yolo.update(dnn_input);
int num_detected = yolo.detected.size();
if (num_detected > MAX_DETECT_SIZE)
num_detected = MAX_DETECT_SIZE;
coords.clear();
end_t = std::chrono::steady_clock::now();
std::cout << " TIME 1 : " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
step_t = end_t;
// draw dets
std::cout << "camera: " << camera->CAM_IDX << " - num detected: " << num_detected << std::endl;
//TODO: move in a thread
// //preprocessing frame
// step_t_segmentation = std::chrono::steady_clock::now();
// // src_gray
// canny_img = img_laplacian(orig_frame,0);
// cv::Canny(canny_img, canny, 100, 100*2 );
// // sprintf(buf_frame_crop_name,"../demo/demo/data/img_disparity/%d_%d_canny.jpg",frame_nbr, 999);
// // cv::imwrite(buf_frame_crop_name, canny);
// end_t_segmentation = std::chrono::steady_clock::now();
// std::cout << " - TIME END pre canny : "<<std::chrono::duration_cast<std::chrono::milliseconds>(end_t_segmentation - step_t_segmentation).count() << " ms"<<std::endl;
// step_t_segmentation = end_t_segmentation;
// // std::cout<<"o: "<<orig_frame.cols<<" - "<<orig_frame.rows<<std::endl;
// // std::cout<<"canny: "<<canny.cols<<" - "<<canny.rows<<std::endl;
// // std::cout<<"pre: "<<pre_canny.cols<<" - "<<pre_canny.rows<<std::endl;
// if(!first_iteration)
// {
// // backtorgb = cv::cvtColor(pre_canny,cv::COLOR_GRAY2RGB)
// cv::cvtColor(pre_canny, pre_canny_RGB, cv::COLOR_GRAY2RGB);
// cv::cvtColor(canny, canny_RGB, cv::COLOR_GRAY2RGB);
// disparity_frame = frame_disparity(pre_canny_RGB, canny_RGB, frame_nbr, 999, 0);
// std::cout<<"size: "<<disparity_frame.rows<<" - "<<disparity_frame.cols<<std::endl;
// if (disparity_frame.rows == 0 || disparity_frame.cols == 0)
// return -1;
// if (disparity_frame.empty())
// { // only fools don't check...
// std::cout << "image not loaded !" << std::endl;
// return -1;
// }
// end_t_segmentation = std::chrono::steady_clock::now();
// std::cout << " TIME canny : frame_disparity : "<<std::chrono::duration_cast<std::chrono::milliseconds>(end_t_segmentation - step_t_segmentation).count() << " ms"<<std::endl;
// step_t_segmentation = end_t_segmentation;
// // //--------------------------------
// // //frame box disparity on the original image
// // step_t_segmentation = std::chrono::steady_clock::now();
// // frame_box_disparity(pre_frame, frame, pre_rois, frame_nbr);
// // // reset pre_rois for the new roi of the current frame
// // // pre_rois.erase(pre_rois.begin(), pre_rois.end());
// // end_t_segmentation = std::chrono::steady_clock::now();
// // std::cout << " TIME Frame disparity : "<<std::chrono::duration_cast<std::chrono::milliseconds>(end_t_segmentation - step_t_segmentation).count() << " ms"<<std::endl;
// // step_t_segmentation = end_t_segmentation;
// // //frame box disparity on the preprocessed image
// // cv::cvtColor(pre_canny, pre_canny_RGB, cv::COLOR_GRAY2RGB);
// // cv::cvtColor(canny, canny_RGB, cv::COLOR_GRAY2RGB);
// // frame_box_disparity(pre_canny_RGB, canny_RGB, pre_rois, frame_nbr);
// // // reset pre_rois for the new roi of the current frame
// // pre_rois.erase(pre_rois.begin(), pre_rois.end());
// // end_t_segmentation = std::chrono::steady_clock::now();
// // std::cout << " TIME Canny Frame disparity : "<<std::chrono::duration_cast<std::chrono::milliseconds>(end_t_segmentation - step_t_segmentation).count() << " ms"<<std::endl;
// // step_t_segmentation = end_t_segmentation;
// // //---------------------------------
// }
// compute some metrics on the whole frame
// segmentation(pre_frame, frame, frame_nbr, 0, 0);
for (int i = 0; i < num_detected; i++)
{
b = yolo.detected[i];
x0 = b.x;
// w = b.w;
// x1 = b.x + w;
// y0 = b.y;
h = b.h;
y1 = b.y + h;
objClass = b.cl;
det_class = obj_class[b.cl];
// prob = b.prob;
intensity = mask.at<uchar>(cv::Point(int(x0 + b.w / 2), y1));
if (intensity[0])
{
if (objClass < 6)
{
// find the rectangular on the frame (sub-figure)
// roi.x = (x0 > 0)? x0 : 0;
// roi.y = (y0 > 0)? y0 : 0;
// // std::cout<<"x "<<roi.x<<" - y "<<roi.y<<std::endl;
// roi.width = (roi.x+w >= frame.cols)? frame.cols-1-roi.x : w;
// roi.height = (roi.y+h >= frame.rows)? frame.rows-1-roi.y : h;
// std::cout<<"w "<<roi.width<<" - h "<<roi.height<<std::endl;
// std::cout<<"wf "<<frame.cols<<" - hf "<<frame.rows<<std::endl;
// std::cout<<"---"<<std::endl;
// std::cout<<"x "<<roi.x<<" to "<<roi.width+roi.x<<" wf "<<frame.cols<<std::endl;
// std::cout<<"y "<<roi.y<<" to "<<roi.height+roi.y<<" hf "<<frame.rows<<std::endl;
//update pre_roi for the next frame
// pre_rois.push_back(roi);
// segmentation(frame(roi), frame(roi), frame_nbr, i, 1);
/////
convert_coords(coords, x0 + b.w / 2, y1, objClass, H, adfGeoTransform);
// //std::cout<<objClass<<" ("<<prob<<"): "<<x0<<" "<<y0<<" "<<x1<<" "<<y1<<"\n";
// cv::rectangle(frame, cv::Point(x0, y0), cv::Point(x1, y1), camera->yolo.colors[objClass], 2);
// // draw label
// int baseline = 0;
// float fontScale = 0.5;
// int thickness = 2;
// cv::Size textSize = getTextSize(det_class, cv::FONT_HERSHEY_SIMPLEX, fontScale, thickness, &baseline);
// cv::rectangle(frame, cv::Point(x0, y0), cv::Point((x0 + textSize.width - 2), (y0 - textSize.height - 2)), camera->yolo.colors[b.cl], -1);
// cv::putText(frame, det_class, cv::Point(x0, (y0 - (baseline / 2))), cv::FONT_HERSHEY_SIMPLEX, fontScale, cv::Scalar(255, 255, 255), thickness);
}
}
}
end_t = std::chrono::steady_clock::now();
std::cout << " TIME 2 : " << std::chrono::duration_cast<std::chrono::milliseconds>(end_t - step_t).count() << " ms" << std::endl;
step_t = end_t;
//convert from latitude and longitude to meters for ekf
cur_frame.clear();
for (size_t i = 0; i < coords.size(); i++)
{
gc.geodetic2Enu(coords[i].lat_, coords[i].long_, 0, &east, &north, &up);
cur_frame.push_back(Data(east, north, frame_nbr, coords[i].class_));
}
if (first_iteration)
{
// if there aren't detections and it is the first iteration, we can't initialize the tracker, so continue
if (cur_frame.empty())
continue;
for (auto f : cur_frame)
trackers.push_back(Tracker(f, initial_age, dt, n_states));
}
else
{
Track(cur_frame, dt, n_states, initial_age, age_threshold, trackers);
}
std::cout << "There are " << trackers.size() << " trackers" << std::endl;
//prepare message with tracker info
if (trackers.size() != 0)
{
// mutex_cv.lock();
addRoadUserfromTracker(trackers, m, gc, maskOrient, adfGeoTransform, H);
// mutex_cv.unlock();
//prepare the message with detection info
//prepare_message(m, coords, CAM_IDX);
//send message
if (!m->objects.empty())
Comm.send_message(m);
}
if (to_show)
{
//populate the ModFrame_t
info_show.sem.lock();
info_show.original_frame = frame.clone();
// std::vector<Tracker> trackers;
info_show.trackers = trackers;
// geodetic_converter::GeodeticConverter gc;
info_show.gc = gc;
for (int i = 0; i < 6; i++)
info_show.adfGeoTransform[i] = adfGeoTransform[i];
// cv::Mat H;
info_show.H = H.clone();
info_show.yolo = yolo;
// std::copy(camera->yolo.begin(), camera->yolo.end(), info_show.yolo.begin());
info_show.mask = mask.clone();
info_show.sem.unlock();
}
// update pre_frame for the disparity map
// pre_frame = orig_frame.clone();
// pre_canny = canny.clone();
if (first_iteration)
first_iteration = false;
frame_nbr++;
std::cout << camera->CAM_IDX << " camera thread: ";
TIMER_STOP
}
return (void *)0;
}
int main(int argc, char *argv[])
{
std::cout << "detection\n";
signal(SIGINT, sig_handler);
srand(time(NULL));
Parameters_t par;
if(!read_parameters(argc, argv, &par))
return -1;
tk::dnn::Yolo3Detection yolo[par.n_cameras];
for(int i=0; i<par.n_cameras; i++)
{
yolo[i].init(par.net);
yolo[i].thresh = 0.25;
// if(SAVE_RESULT)
// resultVideo << frame;
}
// tk::dnn::Yolo3Detection yolo;
// yolo.init(net);
// yolo.thresh = 0.25;
gRun = true;
// start the local clock. It is used to check the incoming frames (by different cameras)
local_clock_start = std::chrono::steady_clock::now();
/*GPS information*/
double *adfGeoTransform = (double *)malloc(6 * sizeof(double));
readTiff(par.tiffile, adfGeoTransform);
// Camera_t cameras[par.n_cameras];
for(int i=0; i<par.n_cameras; i++)
{
for(int j = 0; j < 6; j++ )
par.cameras[i].adfGeoTransform[j] = adfGeoTransform[j];
par.cameras[i].yolo = &yolo[i];
// par.cameras[i].yolo.init(par.net);
// par.cameras[i].yolo.thresh = 0.25;
// cameras[i].yolo = yolo[i];
// cameras[i].yolo = yolo;
}
pthread_t camera_task[par.n_cameras];
for(int i=0; i<par.n_cameras; i++)
{
std::cout<<"creating thread\n";
if(pthread_create(&camera_task[i], NULL, computationTask, (void*)&(par.cameras[i])))
{
fprintf(stderr, "error creating thread\n");
return 1;
}
}
for(int i=0; i<par.n_cameras; i++)
{
pthread_join(camera_task[i], NULL);
}
std::cout <<" free adfGeoT \n";
free(adfGeoTransform);
std::cout << "detection end\n";
return 0;
}
Binary file not shown.
-248
View File
@@ -1,248 +0,0 @@
#ifndef LAYER_H
#define LAYER_H
#include<iostream>
#include "utils.h"
#include "Network.h"
namespace tkDNN {
/**
Data rapresentation beetween layers
n = batch size
c = channels
h = heigth (lines)
w = width (rows)
l = lenght (3rd dimension)
*/
struct dataDim_t {
int n, c, h, w, l;
dataDim_t() : n(1), c(1), h(1), w(1), l(1) {};
dataDim_t(int _n, int _c, int _h, int _w, int _l = 1) :
n(_n), c(_c), h(_h), w(_w), l(_l) {};
void print() {
std::cout<<"Data dim: "<<n<<" "<<c<<" "<<h<<" "<<w<<" "<<l<<"\n";
}
int tot() {
return n*c*h*w*l;
}
};
/**
Simple layer Father class
*/
class Layer {
public:
Layer(Network *net, dataDim_t input_dim);
virtual ~Layer();
virtual value_type* infer(dataDim_t &dim, value_type* srcData) {
std::cout<<"No infer action for this layer\n";
return NULL;
}
dataDim_t input_dim, output_dim;
protected:
Network *net;
cudnnTensorDescriptor_t srcTensorDesc, dstTensorDesc;
};
/**
Father class of all layer that need to load trained weights
*/
class LayerWgs : public Layer {
public:
LayerWgs(Network *net, dataDim_t input_dim,
int inputs, int outputs, int kh, int kw, int kt,
const char* fname_weights, const char* fname_bias);
virtual ~LayerWgs();
protected:
int inputs, outputs;
std::string weights_path, bias_path;
value_type *data_h, *data_d;
value_type *bias_h, *bias_d;
};
/**
Dense (full interconnection) layer
*/
class Dense : public LayerWgs {
public:
Dense(Network *net, dataDim_t in_dim, int out_ch,
const char* fname_weights, const char* fname_bias);
virtual ~Dense();
virtual value_type* infer(dataDim_t &dim, value_type* srcData);
protected:
value_type *dstData; //where results will be putted
};
/**
Avaible activation functions
*/
typedef enum {
ACTIVATION_SIGMOID = 0,
ACTIVATION_RELU = 1,
ACTIVATION_TANH = 2,
ACTIVATION_ELU = 100
} tkdnnActivationMode_t;
/**
Activation layer (it doesnt need weigths)
*/
class Activation : public Layer {
public:
Activation(Network *net, dataDim_t input_dim, tkdnnActivationMode_t act_mode);
virtual ~Activation();
virtual value_type* infer(dataDim_t &dim, value_type* srcData);
protected:
tkdnnActivationMode_t act_mode;
value_type *dstData; //where results will be putted
};
/**
Convolutional 2D layer
*/
class Conv2d : public LayerWgs {
public:
Conv2d(Network *net, dataDim_t in_dim, int out_ch,
int kernelH, int kernelW, int strideH, int strideW,
const char* fname_weights, const char* fname_bias);
virtual ~Conv2d();
virtual value_type* infer(dataDim_t &dim, value_type* srcData);
protected:
value_type *dstData; //where results will be putted
int kernelH, kernelW, strideH, strideW;
cudnnFilterDescriptor_t filterDesc;
cudnnConvolutionDescriptor_t convDesc;
cudnnConvolutionFwdAlgo_t algo;
cudnnTensorDescriptor_t biasTensorDesc;
void* workSpace;
size_t ws_sizeInBytes;
};
/**
Convolutional 3D layer
*/
class Conv3d : public LayerWgs {
public:
Conv3d(Network *net, dataDim_t in_dim, int out_ch,
int kernelH, int kernelW, int kernelL,
int strideH, int strideW, int strideL,
const char* fname_weights, const char* fname_bias);
virtual ~Conv3d();
virtual value_type* infer(dataDim_t &dim, value_type* srcData);
protected:
value_type *dstData; //where results will be putted
int kernelH, kernelW, kernelL;
int strideH, strideW, strideL;
cudnnFilterDescriptor_t filterDesc;
cudnnConvolutionDescriptor_t convDesc;
cudnnConvolutionFwdAlgo_t algo;
cudnnTensorDescriptor_t biasTensorDesc;
cudnnTensorDescriptor_t biasDstTensorDesc;
void* workSpace;
size_t ws_sizeInBytes;
};
/**
Flatten layer
is actually a matrix transposition
*/
class Flatten : public Layer {
public:
Flatten(Network *net, dataDim_t input_dim);
virtual ~Flatten();
virtual value_type* infer(dataDim_t &dim, value_type* srcData);
protected:
value_type *dstData; //where results will be putted
};
/**
MulAdd layer
apply a multiplication and then an addition for each data
*/
class MulAdd : public Layer {
public:
MulAdd(Network *net, dataDim_t input_dim, value_type mul, value_type add);
virtual ~MulAdd();
virtual value_type* infer(dataDim_t &dim, value_type* srcData);
protected:
value_type mul, add;
value_type *dstData, *add_vector; //where results will be putted
};
/**
Avaible pooling functions (padding on tkDNN is not supported)
*/
typedef enum {
POOLING_MAX = 0,
POOLING_AVERAGE = 1, // count for average includes padded values
POOLING_AVERAGE_EXCLUDE_PADDING = 2 // count for average does not include padded values
} tkdnnPoolingMode_t;
/**
Pooling layer
currenty supported only 2d pooing (also on 3d input)
*/
class Pooling : public Layer {
public:
Pooling(Network *net, dataDim_t input_dim, int winH, int winW,
int strideH, int strideW, tkdnnPoolingMode_t pool_mode);
virtual ~Pooling();
virtual value_type* infer(dataDim_t &dim, value_type* srcData);
protected:
cudnnPoolingDescriptor_t poolingDesc;
int winH, winW;
int strideH, strideW;
tkdnnPoolingMode_t pool_mode;
value_type *dstData, *tmpInputData, *tmpOutputData; //where results will be putted
bool poolOn3d;
};
}
#endif //LAYER_H
-36
View File
@@ -1,36 +0,0 @@
#ifndef NETWORK_H
#define NETWORK_H
#include "utils.h"
namespace tkDNN {
struct dataDim_t;
class Layer;
const int MAX_LAYERS = 256;
class Network {
public:
Network();
virtual ~Network();
/**
Do inferece for every added layer
*/
value_type* infer(dataDim_t &dim, value_type* data);
bool addLayer(Layer *l);
cudnnDataType_t dataType;
cudnnTensorFormat_t tensorFormat;
cudnnHandle_t cudnnHandle;
cublasHandle_t cublasHandle;
private:
Layer* layers[MAX_LAYERS]; //contains layers of the net
int num_layers; //current number of layers
};
}
#endif //NETWORK_H
+44
View File
@@ -0,0 +1,44 @@
#include <iostream>
#include <sstream>
#include <fstream>
#include <iomanip>
#include <stdlib.h>
#include <cstring>
#include <cstdlib>
#include <time.h>
#include <chrono>
#include "cuda.h"
#include "cuda_runtime_api.h"
#include <cublas_v2.h>
#include <cudnn.h>
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp>
//saliency
#include <opencv2/core/utility.hpp>
//#include <opencv2/saliency.hpp>
#include <opencv2/highgui.hpp>
#define SAVE false
#define SAVE_TO(name, fn, i, var) {sprintf(buf_frame_crop_name,name,fn,i);\
cv::imwrite(buf_frame_crop_name, var);}
// cv::Mat img_threshold(cv::Mat frame_crop);
// cv::Mat img_background(cv::Mat frame_crop);
// cv::Mat img_dist_transform(cv::Mat frame_crop);
// cv::Mat img_watershed(cv::Mat frame_crop);
void image_segmentation(cv::Mat frame_crop, int frame_nbr, int i);
void image_gradients(cv::Mat frame_crop, int frame_nbr, int i);
void image_find_contours(cv::Mat frame_crop, int frame_nbr, int i);
void image_saliency(cv::Mat frame_crop, int frame_nbr, int i);
void frame_box_disparity(cv::Mat pre_frame, cv::Mat frame, std::vector <cv::Rect> pre_rois, int frame_nbr);
void segmentation(cv::Mat pre_frame, cv::Mat frame_crop, int frame_nbr, int i, int mode);
//canny
cv::Mat img_laplacian(cv::Mat frame_crop, int ret);
cv::Mat frame_disparity(cv::Mat pre_frame, cv::Mat frame, int frame_nbr, int i, int ret);
+32
View File
@@ -0,0 +1,32 @@
#ifndef CALIBRATION_H
#define CALIBRATION_H
#include "gdal.h"
#include <gdal_priv.h>
#include <gdal/gdal.h>
#include "gdal/gdal_priv.h"
#include "gdal/cpl_conv.h"
#include <yaml-cpp/yaml.h>
#include <opencv2/calib3d.hpp>
#include <opencv2/core.hpp>
#include <iostream>
#include <cstring>
struct ObjCoords
{
double lat_;
double long_;
int class_;
};
void readTiff(char *filename, double *adfGeoTransform);
void readCameraCalibrationYaml(const std::string &cameraCalib, cv::Mat &cameraMat, cv::Mat &distCoeff);
void pixel2coord(int x, int y, double &lat, double &lon, double *adfGeoTransform);
void coord2pixel(double lat, double lon, int &x, int &y, double *adfGeoTransform);
void fillMatrix(cv::Mat &H, double *matrix, bool show = false);
void read_projection_matrix(cv::Mat &H, char *path);
void convert_coords(std::vector<ObjCoords> &coords, int x, int y, int detected_class, cv::Mat H, double *adfGeoTransform);
#endif /*CALIBRATION_H*/
+45
View File
@@ -0,0 +1,45 @@
#ifndef CAMERAUTILS_H
#define CAMERAUTILS_H
#include <vector>
#include <mutex>
#include <opencv2/core/core.hpp>
#include "tracker.h"
#include "Yolo3Detection.h"
struct Camera_t
{
int CAM_IDX;
char *input;
char *pmatrix;
char *maskfile;
char *cameraCalib;
char *maskFileOrient;
bool to_show;
tk::dnn::Yolo3Detection *yolo;
double adfGeoTransform[6];
};
struct Frame_t
{
char *input;
cv::Mat frame;
int frame_nbr;
// sem_vc for mainthread, videocapturethread, originalthread and disparitythread
std::mutex sem_vc;
};
struct ModFrame_t
{
std::vector<Tracker> trackers;
geodetic_converter::GeodeticConverter gc;
double adfGeoTransform[6];
cv::Mat H;
cv::Mat original_frame;
tk::dnn::Yolo3Detection yolo;
cv::Mat mask;
// sem for mainthread, detectionthread and topviewthread
std::mutex sem;
};
#endif /*CAMERAUTILS_H*/
+21
View File
@@ -0,0 +1,21 @@
#ifndef CONFIGURATION_H
#define CONFIGURATION_H
#include "cameraUtils.h"
#include <iostream>
#include <cstring>
#include <yaml-cpp/yaml.h>
struct Parameters_t
{
char *net;
char *tiffile;
int n_cameras;
Camera_t *cameras;
};
void readCamerasParametersYaml(const std::string &camerasParams, Parameters_t *par);
bool read_parameters(int argc, char *argv[], Parameters_t *par);
#endif /*CONFIGURATION_H*/
-3
View File
@@ -1,3 +0,0 @@
#include "utils.h"
void activationELUForward(value_type* srcData, value_type* dstData, int size);
+22
View File
@@ -0,0 +1,22 @@
#ifndef MESSAGE_H
#define MESSAGE_H
#include <iostream>
#include <cstdlib>
#include <ctime>
#include <opencv2/calib3d.hpp>
#include <opencv2/core.hpp>
// #include <sys/socket.h> //socket
// #include <arpa/inet.h> //inet_addr
// #include <unistd.h> //write
#include "tracker.h"
#include "../masa_protocol/include/send.hpp"
#include "../masa_protocol/include/serialize.hpp"
unsigned long long time_in_ms();
void addRoadUserfromTracker(const std::vector<Tracker> &trackers, Message *m, geodetic_converter::GeodeticConverter &gc, const cv::Mat &maskOrient, double *adfGeoTransform, cv::Mat H);
#endif /*MESSAGE_H*/
+538
View File
@@ -0,0 +1,538 @@
#ifndef LAYER_H
#define LAYER_H
#include<iostream>
#include<vector>
#include "utils.h"
#include "Network.h"
namespace tk
{
namespace dnn
{
enum layerType_t {
LAYER_INPUT,
LAYER_DENSE,
LAYER_CONV2D,
LAYER_LSTM,
LAYER_ACTIVATION,
LAYER_FLATTEN,
LAYER_MULADD,
LAYER_POOLING,
LAYER_SOFTMAX,
LAYER_ROUTE,
LAYER_REORG,
LAYER_SHORTCUT,
LAYER_UPSAMPLE,
LAYER_REGION,
LAYER_YOLO
};
#define TKDNN_BN_MIN_EPSILON 1e-5
/**
Simple layer Father class
*/
class Layer
{
public:
Layer(Network *net);
virtual ~Layer();
virtual layerType_t getLayerType() = 0;
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData)
{
std::cout << "No infer action for this layer\n";
return NULL;
}
dataDim_t input_dim, output_dim;
dnnType *dstData; //where results will be putted
std::string getLayerName()
{
layerType_t type = getLayerType();
switch(type) {
case LAYER_INPUT: return "Input";
case LAYER_DENSE: return "Dense";
case LAYER_CONV2D: return "Conv2d";
case LAYER_LSTM: return "LSTM";
case LAYER_ACTIVATION: return "Activation";
case LAYER_FLATTEN: return "Flatten";
case LAYER_MULADD: return "MulAdd";
case LAYER_POOLING: return "Pooling";
case LAYER_SOFTMAX: return "Softmax";
case LAYER_ROUTE: return "Route";
case LAYER_REORG: return "Reorg";
case LAYER_SHORTCUT: return "Shortcut";
case LAYER_UPSAMPLE: return "Upsample";
case LAYER_REGION: return "Region";
case LAYER_YOLO: return "Yolo";
default: return "unknown";
}
}
protected:
Network *net;
cudnnTensorDescriptor_t srcTensorDesc, dstTensorDesc;
};
/**
Father class of all layer that need to load trained weights
*/
class LayerWgs : public Layer
{
public:
LayerWgs(Network *net, int inputs, int outputs, int kh, int kw, int kt,
std::string fname_weights, bool batchnorm = false);
virtual ~LayerWgs();
int inputs, outputs;
std::string weights_path;
dnnType *data_h, *data_d;
dnnType *bias_h, *bias_d;
//batchnorm
bool batchnorm;
dnnType *power_h;
dnnType *scales_h, *scales_d;
dnnType *mean_h, *mean_d;
dnnType *variance_h, *variance_d;
//fp16
__half *data16_h, *bias16_h;
__half *data16_d, *bias16_d;
__half *power16_h, *power16_d;
__half *scales16_h, *scales16_d;
__half *mean16_h, *mean16_d;
__half *variance16_h, *variance16_d;
};
/**
Input layer (it doesnt need weigths)
*/
class Input : public Layer {
public:
Input(Network *net, dataDim_t &dim, dnnType* srcData) : Layer(net) {
input_dim = dim;
output_dim = dim;
dstData = srcData;
}
virtual ~Input() {}
virtual layerType_t getLayerType() { return LAYER_INPUT; };
virtual dnnType* infer(dataDim_t &dim, dnnType* srcData) {
return dstData;
}
};
/**
Dense (full interconnection) layer
*/
class Dense : public LayerWgs
{
public:
Dense(Network *net, int out_ch, std::string fname_weights);
virtual ~Dense();
virtual layerType_t getLayerType() { return LAYER_DENSE; };
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData);
};
/**
Avaible activation functions
*/
typedef enum
{
ACTIVATION_ELU = 100,
ACTIVATION_LEAKY = 101
} tkdnnActivationMode_t;
/**
Activation layer (it doesnt need weigths)
*/
class Activation : public Layer
{
public:
int act_mode;
Activation(Network *net, int act_mode);
virtual ~Activation();
virtual layerType_t getLayerType() { return LAYER_ACTIVATION; };
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData);
protected:
cudnnActivationDescriptor_t activDesc;
};
/**
Convolutional 2D layer
WEIGHTS shape: OUTCH, INCH, KH, KW ...
BIAS shape: OUTCH
with BATCHNORM:
scales: OUTCH
means: OUTCH
variance: OUTCH
*/
class Conv2d : public LayerWgs
{
public:
Conv2d( Network *net, int out_ch, int kernelH, int kernelW,
int strideH, int strideW, int paddingH, int paddingW,
std::string fname_weights, bool batchnorm = false);
virtual ~Conv2d();
virtual layerType_t getLayerType() { return LAYER_CONV2D; };
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData);
int kernelH, kernelW, strideH, strideW, paddingH, paddingW;
protected:
cudnnFilterDescriptor_t filterDesc;
cudnnConvolutionDescriptor_t convDesc;
cudnnConvolutionFwdAlgo_t algo;
cudnnTensorDescriptor_t biasTensorDesc;
void *workSpace;
size_t ws_sizeInBytes;
};
/**
Bidirectional LSTM layer
ONLY BIDIRECTIONAL (TODO: more configurable)
currently implemented as 2 inferences: forward and backward (TODO: only 1 cudnn inference)
implementation info:
https://github.com/jiangnanhugo/seq2seq_cuda/blob/e4dbdcfa0517c972bfd4beea9f11a5233954093c/src/rnn.cpp
https://github.com/Jeffery-Song/mxnet-test/blob/aab666faad44011f7a67b527b5f6c960367d0422/src/operator/cudnn_rnn-inl.h
https://stackoverflow.com/a/38737941
https://colah.github.io/posts/2015-08-Understanding-LSTMs/
PARAMS (numlayers*2):
layer0:
( INCH, ? ) ???
( HIDDEN, ? ) ???
( HIDDEN * 8 ) ???
layer2:
( INCH, ? ) ???
( HIDDEN, ? ) ???
( HIDDEN * 8 ) ???
OUTPUT shape:
(N, C, 1, W) ---> LSTM(HIDDEN, returnSeq=True) ---> (N, 2*HIDDEN, 1, W) # W is seqLength
(N, C, 1, W) ---> LSTM(HIDDEN, returnSeq=False) ---> (N, 2*HIDDEN, 1, 1)
*/
class LSTM : public Layer {
public:
LSTM(Network *net, int hiddensize, bool returnSeq, std::string fname_weights);
virtual ~LSTM();
virtual layerType_t getLayerType() { return LAYER_LSTM; };
virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
const bool bidirectional = true; /**> is the net bidir */
bool returnSeq = false; /**> if false return only the result of last timestep */
int stateSize = 0; /**> number of hidden states */
int seqLen = 0; /**> number of timesteps */
int numLayers = 1; /**> number of internal layers */
protected:
cudnnRNNDescriptor_t rnnDesc;
cudnnDropoutDescriptor_t dropoutDesc;
dnnType *dropout_states_, *work_space_;
size_t workspace_byte_, dropout_byte_;
int workspace_size_, dropout_size_;
std::vector<cudnnTensorDescriptor_t> x_desc_vec_, y_desc_vec_;
cudnnTensorDescriptor_t hx_desc_, cx_desc_;
cudnnTensorDescriptor_t hy_desc_, cy_desc_;
dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr;
int stateDataDim;
cudnnFilterDescriptor_t w_desc_;
dnnType *w_ptr;
dnnType *w_h;
dnnType *wf_ptr, *wb_ptr; // params pointer forward and backward layer
// used during inference
dataDim_t one_output_dim; // output dim of as single inference
dnnType *srcF, *srcB; // input of single inference
dnnType *dstF, *dstB_NR, *dstB; // output of single inference, dstB_NR = dstB not reversed
};
/**
Flatten layer
is actually a matrix transposition
*/
class Flatten : public Layer
{
public:
Flatten(Network *net);
virtual ~Flatten();
virtual layerType_t getLayerType() { return LAYER_FLATTEN; };
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData);
};
/**
MulAdd layer
apply a multiplication and then an addition for each data
*/
class MulAdd : public Layer
{
public:
MulAdd(Network *net, dnnType mul, dnnType add);
virtual ~MulAdd();
virtual layerType_t getLayerType() { return LAYER_MULADD; };
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData);
protected:
dnnType mul, add;
dnnType *add_vector;
};
/**
Avaible pooling functions (padding on tkDNN is not supported)
*/
typedef enum
{
POOLING_MAX = 0,
POOLING_AVERAGE = 1, // count for average includes padded values
POOLING_AVERAGE_EXCLUDE_PADDING = 2 // count for average does not include padded values
} tkdnnPoolingMode_t;
/**
Pooling layer
currenty supported only 2d pooing (also on 3d input)
*/
class Pooling : public Layer
{
public:
int winH, winW;
int strideH, strideW;
int paddingH, paddingW;
Pooling(Network *net, int winH, int winW,
int strideH, int strideW, tkdnnPoolingMode_t pool_mode);
virtual ~Pooling();
virtual layerType_t getLayerType() { return LAYER_POOLING; };
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData);
protected:
cudnnPoolingDescriptor_t poolingDesc;
tkdnnPoolingMode_t pool_mode;
dnnType *tmpInputData, *tmpOutputData;
bool poolOn3d;
};
/**
Softmax layer
*/
class Softmax : public Layer
{
public:
Softmax(Network *net);
virtual ~Softmax();
virtual layerType_t getLayerType() { return LAYER_SOFTMAX; };
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData);
};
/**
Route layer
Merge a list of layers
*/
class Route : public Layer
{
public:
Route(Network *net, Layer **layers, int layers_n);
virtual ~Route();
virtual layerType_t getLayerType() { return LAYER_ROUTE; };
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData);
public:
Layer **layers; //ids of layers to be merged
int layers_n; //number of layers
};
/**
Reorg layer
Mantain same dimension but change C*H*W distribution
*/
class Reorg : public Layer
{
public:
Reorg(Network *net, int stride);
virtual ~Reorg();
virtual layerType_t getLayerType() { return LAYER_REORG; };
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData);
int stride;
};
/**
Shortcut layer
sum with stride another layer
*/
class Shortcut : public Layer
{
public:
Shortcut(Network *net, Layer *backLayer);
virtual ~Shortcut();
virtual layerType_t getLayerType() { return LAYER_SHORTCUT; };
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData);
public:
Layer *backLayer;
};
/**
Upsample layer
Mantain same dimension but change C*H*W distribution
*/
class Upsample : public Layer
{
public:
Upsample(Network *net, int stride);
virtual ~Upsample();
virtual layerType_t getLayerType() { return LAYER_UPSAMPLE; };
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData);
int stride;
bool reverse;
};
struct box
{
int cl;
float x, y, w, h;
float prob;
};
struct sortable_bbox
{
int index;
int cl;
float **probs;
};
/**
Yolo3 layer
*/
class Yolo : public Layer
{
public:
struct box
{
float x, y, w, h;
};
struct detection
{
Yolo::box bbox;
int classes;
float *prob;
float *mask;
float objectness;
int sort_class;
};
Yolo(Network *net, int classes, int num, std::string fname_weights);
virtual ~Yolo();
virtual layerType_t getLayerType() { return LAYER_YOLO; };
int classes, num;
dnnType *mask_h, *mask_d; //anchors
dnnType *bias_h, *bias_d; //anchors
std::vector<std::string> classesNames;
virtual dnnType *infer(dataDim_t &dim, dnnType *srcData);
int computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh);
dnnType *predictions;
static const int MAX_DETECTIONS = 256;
static Yolo::detection *allocateDetections(int nboxes, int classes);
static void mergeDetections(Yolo::detection *dets, int ndets, int classes);
};
/**
Region layer
*/
class Region : public Layer
{
public:
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);
};
class RegionInterpret
{
public:
RegionInterpret(dataDim_t input_dim, dataDim_t output_dim,
int classes, int coords, int num, float thresh, std::string 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;
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(dnnType *data_h, int imageW = 0, int imageH = 0);
void showImageResult(dnnType *input_h);
static float box_iou(box a, box b);
};
} // namespace dnn
} // namespace tk
#endif //LAYER_H
+73
View File
@@ -0,0 +1,73 @@
#ifndef NETWORK_H
#define NETWORK_H
#include "utils.h"
namespace tk
{
namespace dnn
{
/**
Data rapresentation beetween layers
n = batch size
c = channels
h = heigth (lines)
w = width (rows)
l = lenght (3rd dimension)
*/
struct dataDim_t
{
int n, c, h, w, l;
dataDim_t() : n(1), c(1), h(1), w(1), l(1){};
dataDim_t(int _n, int _c, int _h, int _w, int _l = 1) : n(_n), c(_c), h(_h), w(_w), l(_l){};
void print()
{
std::cout << "Data dim: " << n << " " << c << " " << h << " " << w << " " << l << "\n";
}
int tot()
{
return n * c * h * w * l;
}
};
class Layer;
const int MAX_LAYERS = 256;
class Network
{
public:
Network(dataDim_t input_dim);
virtual ~Network();
/**
Do inferece for every added layer
*/
dnnType *infer(dataDim_t &dim, dnnType *data);
bool addLayer(Layer *l);
void print();
cudnnDataType_t dataType;
cudnnTensorFormat_t tensorFormat;
cudnnHandle_t cudnnHandle;
cublasHandle_t cublasHandle;
Layer *layers[MAX_LAYERS]; //contains layers of the net
int num_layers; //current number of layers
dataDim_t input_dim;
dataDim_t getOutputDim();
bool fp16, dla;
};
} // namespace dnn
} // namespace tk
#endif //NETWORK_H
+99
View File
@@ -0,0 +1,99 @@
#ifndef NETWORKRT_H
#define NETWORKRT_H
#include <string.h> // memcpy
#include "utils.h"
#include "Network.h"
#include "Layer.h"
#include "NvInfer.h"
namespace tk
{
namespace dnn
{
template <typename T>
void writeBUF(char *&buffer, const T &val)
{
*reinterpret_cast<T *>(buffer) = val;
buffer += sizeof(T);
}
template <typename T>
T readBUF(const char *&buffer)
{
T val = *reinterpret_cast<const T *>(buffer);
buffer += sizeof(T);
return val;
}
using namespace nvinfer1;
#include "pluginsRT/ActivationLeakyRT.h"
#include "pluginsRT/ReorgRT.h"
#include "pluginsRT/RegionRT.h"
//#include "pluginsRT/RouteRT.h"
#include "pluginsRT/ShortcutRT.h"
#include "pluginsRT/YoloRT.h"
#include "pluginsRT/UpsampleRT.h"
//#include "pluginsRT/Int8Calibrator.h"
class PluginFactory : IPluginFactory
{
public:
YoloRT *yolos[16];
int n_yolos;
virtual IPlugin *createPlugin(const char *layerName, const void *serialData, size_t serialLength);
};
class NetworkRT
{
public:
nvinfer1::DataType dtRT;
nvinfer1::IBuilder *builderRT;
nvinfer1::IRuntime *runtimeRT;
nvinfer1::INetworkDefinition *networkRT;
nvinfer1::ICudaEngine *engineRT;
nvinfer1::IExecutionContext *contextRT;
const static int MAX_BUFFERS_RT = 10;
void *buffersRT[MAX_BUFFERS_RT];
int buf_input_idx, buf_output_idx;
dataDim_t input_dim, output_dim;
dnnType *output;
cudaStream_t stream;
PluginFactory *pluginFactory;
NetworkRT(Network *net, const char *name);
virtual ~NetworkRT();
/**
Do inferece
*/
dnnType *infer(dataDim_t &dim, dnnType *data);
void enqueue();
nvinfer1::ILayer *convert_layer(nvinfer1::ITensor *input, Layer *l);
nvinfer1::ILayer *convert_layer(nvinfer1::ITensor *input, Conv2d *l);
nvinfer1::ILayer *convert_layer(nvinfer1::ITensor *input, Activation *l);
nvinfer1::ILayer *convert_layer(nvinfer1::ITensor *input, Dense *l);
nvinfer1::ILayer *convert_layer(nvinfer1::ITensor *input, Pooling *l);
nvinfer1::ILayer *convert_layer(nvinfer1::ITensor *input, Softmax *l);
nvinfer1::ILayer *convert_layer(nvinfer1::ITensor *input, Route *l);
nvinfer1::ILayer *convert_layer(nvinfer1::ITensor *input, Reorg *l);
nvinfer1::ILayer *convert_layer(nvinfer1::ITensor *input, Region *l);
nvinfer1::ILayer *convert_layer(nvinfer1::ITensor *input, Shortcut *l);
nvinfer1::ILayer *convert_layer(nvinfer1::ITensor *input, Yolo *l);
nvinfer1::ILayer *convert_layer(nvinfer1::ITensor *input, Upsample *l);
bool serialize(const char *filename);
bool deserialize(const char *filename);
};
} // namespace dnn
} // namespace tk
#endif //NETWORKRT_H
+78
View File
@@ -0,0 +1,78 @@
#ifndef YOLO3DDETECTION_H
#define YOLO3DDETECTION_H
#include <iostream>
#include <signal.h>
#include <stdlib.h> /* srand, rand */
#include <unistd.h>
#include <mutex>
#include "utils.h"
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp>
#include "tkdnn.h"
namespace tk
{
namespace dnn
{
/**
*
* @author Francesco Gatti
*/
class Yolo3Detection
{
private:
tk::dnn::NetworkRT *netRT = nullptr;
tk::dnn::Yolo *yolo[3];
dnnType *input, *input_d;
int ndets = 0;
tk::dnn::Yolo::detection *dets = nullptr;
cv::Mat imageF;
cv::Mat bgr[3];
public:
int classes = 0;
int num = 0;
float thresh = 0.3;
cv::Scalar colors[256];
// this is filled with results
std::vector<tk::dnn::box> detected;
// keep track of inference times (ms)
std::vector<double> stats;
Yolo3Detection() {}
virtual ~Yolo3Detection() {}
/**
* Method used to inizialize the class
*
* @return Success of the initialization
*/
bool init(std::string tensor_path);
void addBorders(cv::Mat &imageORIG, cv::Mat &imageWBorders, int &top, int &left);
void update(cv::Mat &frame);
tk::dnn::Yolo* getYoloLayer(int n=0) {
if(n<3)
return yolo[n];
else
return nullptr;
}
};
} // namespace dnn
} // namespace tk
#endif /*YOLO3DDETECTION_H*/
+27
View File
@@ -0,0 +1,27 @@
#ifndef KERNELS_H
#define KERNELS_H
#include "utils.h"
void activationELUForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
void activationLOGISTICForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
void fill(dnnType* data, int size, dnnType val, cudaStream_t stream = cudaStream_t(0));
void reorgForward( dnnType* srcData, dnnType* dstData,
int n, int c, int h, int w, int stride, cudaStream_t stream = cudaStream_t(0));
void softmaxForward(float *input, int n, int batch, int batch_offset,
int groups, int group_offset, int stride, float temp, float *output, cudaStream_t stream = cudaStream_t(0));
void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1, int w1, int s1,
int n2, int c2, int h2, int w2, int s2,
cudaStream_t stream = cudaStream_t(0));
void upsampleForward(dnnType* srcData, dnnType* dstData,
int n, int c, int h, int w, int s, int forward, float scale,
cudaStream_t stream = cudaStream_t(0));
void float2half(float* srcData, __half* dstData, int size, const cudaStream_t stream = cudaStream_t(0));
#endif //KERNELS_H
+289
View File
@@ -0,0 +1,289 @@
int preYoloFilters = (classes+5)*3;
std::string input_bin = bin_path + "/layers/input.bin";
std::vector<std::string> output_bins = {
bin_path + "/debug/layer82_out.bin",
bin_path + "/debug/layer94_out.bin",
bin_path + "/debug/layer106_out.bin"
};
std::string c0_bin = bin_path + "/layers/c0.bin";
std::string c1_bin = bin_path + "/layers/c1.bin";
std::string c2_bin = bin_path + "/layers/c2.bin";
std::string c3_bin = bin_path + "/layers/c3.bin";
std::string c5_bin = bin_path + "/layers/c5.bin";
std::string c6_bin = bin_path + "/layers/c6.bin";
std::string c7_bin = bin_path + "/layers/c7.bin";
std::string c9_bin = bin_path + "/layers/c9.bin";
std::string c10_bin = bin_path + "/layers/c10.bin";
std::string c12_bin = bin_path + "/layers/c12.bin";
std::string c13_bin = bin_path + "/layers/c13.bin";
std::string c14_bin = bin_path + "/layers/c14.bin";
std::string c16_bin = bin_path + "/layers/c16.bin";
std::string c17_bin = bin_path + "/layers/c17.bin";
std::string c19_bin = bin_path + "/layers/c19.bin";
std::string c20_bin = bin_path + "/layers/c20.bin";
std::string c22_bin = bin_path + "/layers/c22.bin";
std::string c23_bin = bin_path + "/layers/c23.bin";
std::string c25_bin = bin_path + "/layers/c25.bin";
std::string c26_bin = bin_path + "/layers/c26.bin";
std::string c28_bin = bin_path + "/layers/c28.bin";
std::string c29_bin = bin_path + "/layers/c29.bin";
std::string c31_bin = bin_path + "/layers/c31.bin";
std::string c32_bin = bin_path + "/layers/c32.bin";
std::string c34_bin = bin_path + "/layers/c34.bin";
std::string c35_bin = bin_path + "/layers/c35.bin";
std::string c37_bin = bin_path + "/layers/c37.bin";
std::string c38_bin = bin_path + "/layers/c38.bin";
std::string c39_bin = bin_path + "/layers/c39.bin";
std::string c41_bin = bin_path + "/layers/c41.bin";
std::string c42_bin = bin_path + "/layers/c42.bin";
std::string c44_bin = bin_path + "/layers/c44.bin";
std::string c45_bin = bin_path + "/layers/c45.bin";
std::string c47_bin = bin_path + "/layers/c47.bin";
std::string c48_bin = bin_path + "/layers/c48.bin";
std::string c50_bin = bin_path + "/layers/c50.bin";
std::string c51_bin = bin_path + "/layers/c51.bin";
std::string c53_bin = bin_path + "/layers/c53.bin";
std::string c54_bin = bin_path + "/layers/c54.bin";
std::string c56_bin = bin_path + "/layers/c56.bin";
std::string c57_bin = bin_path + "/layers/c57.bin";
std::string c59_bin = bin_path + "/layers/c59.bin";
std::string c60_bin = bin_path + "/layers/c60.bin";
std::string c62_bin = bin_path + "/layers/c62.bin";
std::string c63_bin = bin_path + "/layers/c63.bin";
std::string c64_bin = bin_path + "/layers/c64.bin";
std::string c66_bin = bin_path + "/layers/c66.bin";
std::string c67_bin = bin_path + "/layers/c67.bin";
std::string c69_bin = bin_path + "/layers/c69.bin";
std::string c70_bin = bin_path + "/layers/c70.bin";
std::string c72_bin = bin_path + "/layers/c72.bin";
std::string c73_bin = bin_path + "/layers/c73.bin";
std::string c75_bin = bin_path + "/layers/c75.bin";
std::string c76_bin = bin_path + "/layers/c76.bin";
std::string c77_bin = bin_path + "/layers/c77.bin";
std::string c78_bin = bin_path + "/layers/c78.bin";
std::string c79_bin = bin_path + "/layers/c79.bin";
std::string c80_bin = bin_path + "/layers/c80.bin";
std::string c81_bin = bin_path + "/layers/c81.bin";
std::string g82_bin = bin_path + "/layers/g82.bin";
std::string c84_bin = bin_path + "/layers/c84.bin";
std::string c87_bin = bin_path + "/layers/c87.bin";
std::string c88_bin = bin_path + "/layers/c88.bin";
std::string c89_bin = bin_path + "/layers/c89.bin";
std::string c90_bin = bin_path + "/layers/c90.bin";
std::string c91_bin = bin_path + "/layers/c91.bin";
std::string c92_bin = bin_path + "/layers/c92.bin";
std::string c93_bin = bin_path + "/layers/c93.bin";
std::string g94_bin = bin_path + "/layers/g94.bin";
std::string c96_bin = bin_path + "/layers/c96.bin";
std::string c99_bin = bin_path + "/layers/c99.bin";
std::string c100_bin = bin_path + "/layers/c100.bin";
std::string c101_bin = bin_path + "/layers/c101.bin";
std::string c102_bin = bin_path + "/layers/c102.bin";
std::string c103_bin = bin_path + "/layers/c103.bin";
std::string c104_bin = bin_path + "/layers/c104.bin";
std::string c105_bin = bin_path + "/layers/c105.bin";
std::string g106_bin = bin_path + "/layers/g106.bin";
tk::dnn::Conv2d c0 (&net, 32, 3, 3, 1, 1, 1, 1, c0_bin, true);
tk::dnn::Activation a0 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c1 (&net, 64, 3, 3, 2, 2, 1, 1, c1_bin, true);
tk::dnn::Activation a1 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c2 (&net, 32, 1, 1, 1, 1, 0, 0, c2_bin, true);
tk::dnn::Activation a2 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c3 (&net, 64, 3, 3, 1, 1, 1, 1, c3_bin, true);
tk::dnn::Activation a3 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s4 (&net, &a1);
tk::dnn::Conv2d c5 (&net, 128, 3, 3, 2, 2, 1, 1, c5_bin, true);
tk::dnn::Activation a5 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c6 (&net, 64, 1, 1, 1, 1, 0, 0, c6_bin, true);
tk::dnn::Activation a6 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c7 (&net, 128, 3, 3, 1, 1, 1, 1, c7_bin, true);
tk::dnn::Activation a7 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s8 (&net, &a5);
tk::dnn::Conv2d c9 (&net, 64, 1, 1, 1, 1, 0, 0, c9_bin, true);
tk::dnn::Activation a9 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c10 (&net, 128, 3, 3, 1, 1, 1, 1, c10_bin, true);
tk::dnn::Activation a10 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s11 (&net, &s8);
tk::dnn::Conv2d c12 (&net, 256, 3, 3, 2, 2, 1, 1, c12_bin, true);
tk::dnn::Activation a12 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c13 (&net, 128, 1, 1, 1, 1, 0, 0, c13_bin, true);
tk::dnn::Activation a13 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c14 (&net, 256, 3, 3, 1, 1, 1, 1, c14_bin, true);
tk::dnn::Activation a14 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s15 (&net, &a12);
tk::dnn::Conv2d c16 (&net, 128, 1, 1, 1, 1, 0, 0, c16_bin, true);
tk::dnn::Activation a16 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c17 (&net, 256, 3, 3, 1, 1, 1, 1, c17_bin, true);
tk::dnn::Activation a17 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s18 (&net, &s15);
tk::dnn::Conv2d c19 (&net, 128, 1, 1, 1, 1, 0, 0, c19_bin, true);
tk::dnn::Activation a19 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c20 (&net, 256, 3, 3, 1, 1, 1, 1, c20_bin, true);
tk::dnn::Activation a20 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s21 (&net, &s18);
tk::dnn::Conv2d c22 (&net, 128, 1, 1, 1, 1, 0, 0, c22_bin, true);
tk::dnn::Activation a22 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c23 (&net, 256, 3, 3, 1, 1, 1, 1, c23_bin, true);
tk::dnn::Activation a23 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s24 (&net, &s21);
tk::dnn::Conv2d c25 (&net, 128, 1, 1, 1, 1, 0, 0, c25_bin, true);
tk::dnn::Activation a25 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c26 (&net, 256, 3, 3, 1, 1, 1, 1, c26_bin, true);
tk::dnn::Activation a26 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s27 (&net, &s24);
tk::dnn::Conv2d c28 (&net, 128, 1, 1, 1, 1, 0, 0, c28_bin, true);
tk::dnn::Activation a28 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c29 (&net, 256, 3, 3, 1, 1, 1, 1, c29_bin, true);
tk::dnn::Activation a29 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s30 (&net, &s27);
tk::dnn::Conv2d c31 (&net, 128, 1, 1, 1, 1, 0, 0, c31_bin, true);
tk::dnn::Activation a31 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c32 (&net, 256, 3, 3, 1, 1, 1, 1, c32_bin, true);
tk::dnn::Activation a32 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s33 (&net, &s30);
tk::dnn::Conv2d c34 (&net, 128, 1, 1, 1, 1, 0, 0, c34_bin, true);
tk::dnn::Activation a34 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c35 (&net, 256, 3, 3, 1, 1, 1, 1, c35_bin, true);
tk::dnn::Activation a35 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s36 (&net, &s33);
tk::dnn::Conv2d c37 (&net, 512, 3, 3, 2, 2, 1, 1, c37_bin, true);
tk::dnn::Activation a37 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c38 (&net, 256, 1, 1, 1, 1, 0, 0, c38_bin, true);
tk::dnn::Activation a38 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c39 (&net, 512, 3, 3, 1, 1, 1, 1, c39_bin, true);
tk::dnn::Activation a39 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s40 (&net, &a37);
tk::dnn::Conv2d c41 (&net, 256, 1, 1, 1, 1, 0, 0, c41_bin, true);
tk::dnn::Activation a41 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c42 (&net, 512, 3, 3, 1, 1, 1, 1, c42_bin, true);
tk::dnn::Activation a42 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s43 (&net, &s40);
tk::dnn::Conv2d c44 (&net, 256, 1, 1, 1, 1, 0, 0, c44_bin, true);
tk::dnn::Activation a44 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c45 (&net, 512, 3, 3, 1, 1, 1, 1, c45_bin, true);
tk::dnn::Activation a45 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s46 (&net, &s43);
tk::dnn::Conv2d c47 (&net, 256, 1, 1, 1, 1, 0, 0, c47_bin, true);
tk::dnn::Activation a47 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c48 (&net, 512, 3, 3, 1, 1, 1, 1, c48_bin, true);
tk::dnn::Activation a48 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s49 (&net, &s46);
tk::dnn::Conv2d c50 (&net, 256, 1, 1, 1, 1, 0, 0, c50_bin, true);
tk::dnn::Activation a50 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c51 (&net, 512, 3, 3, 1, 1, 1, 1, c51_bin, true);
tk::dnn::Activation a51 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s52 (&net, &s49);
tk::dnn::Conv2d c53 (&net, 256, 1, 1, 1, 1, 0, 0, c53_bin, true);
tk::dnn::Activation a53 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c54 (&net, 512, 3, 3, 1, 1, 1, 1, c54_bin, true);
tk::dnn::Activation a54 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s55 (&net, &s52);
tk::dnn::Conv2d c56 (&net, 256, 1, 1, 1, 1, 0, 0, c56_bin, true);
tk::dnn::Activation a56 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c57 (&net, 512, 3, 3, 1, 1, 1, 1, c57_bin, true);
tk::dnn::Activation a57 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s58 (&net, &s55);
tk::dnn::Conv2d c59 (&net, 256, 1, 1, 1, 1, 0, 0, c59_bin, true);
tk::dnn::Activation a59 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c60 (&net, 512, 3, 3, 1, 1, 1, 1, c60_bin, true);
tk::dnn::Activation a60 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s61 (&net, &s58);
tk::dnn::Conv2d c62 (&net,1024, 3, 3, 2, 2, 1, 1, c62_bin, true);
tk::dnn::Activation a62 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c63 (&net, 512, 1, 1, 1, 1, 0, 0, c63_bin, true);
tk::dnn::Activation a63 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c64 (&net,1024, 3, 3, 1, 1, 1, 1, c64_bin, true);
tk::dnn::Activation a64 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s65 (&net, &a62);
tk::dnn::Conv2d c66 (&net, 512, 1, 1, 1, 1, 0, 0, c66_bin, true);
tk::dnn::Activation a66 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c67 (&net,1024, 3, 3, 1, 1, 1, 1, c67_bin, true);
tk::dnn::Activation a67 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s68 (&net, &s65);
tk::dnn::Conv2d c69 (&net, 512, 1, 1, 1, 1, 0, 0, c69_bin, true);
tk::dnn::Activation a69 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c70 (&net,1024, 3, 3, 1, 1, 1, 1, c70_bin, true);
tk::dnn::Activation a70 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s71 (&net, &s68);
tk::dnn::Conv2d c72 (&net, 512, 1, 1, 1, 1, 0, 0, c72_bin, true);
tk::dnn::Activation a72 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c73 (&net,1024, 3, 3, 1, 1, 1, 1, c73_bin, true);
tk::dnn::Activation a73 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Shortcut s74 (&net, &s71);
tk::dnn::Conv2d c75 (&net, 512, 1, 1, 1, 1, 0, 0, c75_bin, true);
tk::dnn::Activation a75 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c76 (&net,1024, 3, 3, 1, 1, 1, 1, c76_bin, true);
tk::dnn::Activation a76 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c77 (&net, 512, 1, 1, 1, 1, 0, 0, c77_bin, true);
tk::dnn::Activation a77 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c78 (&net,1024, 3, 3, 1, 1, 1, 1, c78_bin, true);
tk::dnn::Activation a78 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c79 (&net, 512, 1, 1, 1, 1, 0, 0, c79_bin, true);
tk::dnn::Activation a79 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c80 (&net,1024, 3, 3, 1, 1, 1, 1, c80_bin, true);
tk::dnn::Activation a80 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c81 (&net, preYoloFilters, 1, 1, 1, 1, 0, 0, c81_bin, false);
tk::dnn::Yolo yolo0 (&net, classes, 3, g82_bin);
tk::dnn::Layer *m83_layers[1] = { &a79 };
tk::dnn::Route m83 (&net, m83_layers, 1);
tk::dnn::Conv2d c84 (&net, 256, 1, 1, 1, 1, 0, 0, c84_bin, true);
tk::dnn::Activation a84 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Upsample u85 (&net, 2);
tk::dnn::Layer *m86_layers[2] = { &u85, &s61 };
tk::dnn::Route m86 (&net, m86_layers, 2);
tk::dnn::Conv2d c87 (&net, 256, 1, 1, 1, 1, 0, 0, c87_bin, true);
tk::dnn::Activation a87 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c88 (&net, 512, 3, 3, 1, 1, 1, 1, c88_bin, true);
tk::dnn::Activation a88 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c89 (&net, 256, 1, 1, 1, 1, 0, 0, c89_bin, true);
tk::dnn::Activation a89 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c90 (&net, 512, 3, 3, 1, 1, 1, 1, c90_bin, true);
tk::dnn::Activation a90 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c91 (&net, 256, 1, 1, 1, 1, 0, 0, c91_bin, true);
tk::dnn::Activation a91 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c92 (&net, 512, 3, 3, 1, 1, 1, 1, c92_bin, true);
tk::dnn::Activation a92 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c93 (&net, preYoloFilters, 1, 1, 1, 1, 0, 0, c93_bin, false);
tk::dnn::Yolo yolo1 (&net, classes, 3, g94_bin);
tk::dnn::Layer *m95_layers[1] = { &a91 };
tk::dnn::Route m95 (&net, m95_layers, 1);
tk::dnn::Conv2d c96 (&net, 128, 1, 1, 1, 1, 0, 0, c96_bin, true);
tk::dnn::Activation a96 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Upsample u97 (&net, 2);
tk::dnn::Layer *m98_layers[2] = { &u97, &s36 };
tk::dnn::Route m98 (&net, m98_layers, 2);
tk::dnn::Conv2d c99 (&net, 128, 1, 1, 1, 1, 0, 0, c99_bin, true);
tk::dnn::Activation a99 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c100 (&net, 256, 3, 3, 1, 1, 1, 1, c100_bin, true);
tk::dnn::Activation a100 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c101 (&net, 128, 1, 1, 1, 1, 0, 0, c101_bin, true);
tk::dnn::Activation a101 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c102 (&net, 256, 3, 3, 1, 1, 1, 1, c102_bin, true);
tk::dnn::Activation a102 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c103 (&net, 128, 1, 1, 1, 1, 0, 0, c103_bin, true);
tk::dnn::Activation a103 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c104 (&net, 256, 3, 3, 1, 1, 1, 1, c104_bin, true);
tk::dnn::Activation a104 (&net, tk::dnn::ACTIVATION_LEAKY);
tk::dnn::Conv2d c105 (&net, preYoloFilters, 1, 1, 1, 1, 0, 0, c105_bin, false);
tk::dnn::Yolo yolo2 (&net, classes, 3, g106_bin);
yolo[0] = &yolo0;
yolo[1] = &yolo1;
yolo[2] = &yolo2;
@@ -0,0 +1,60 @@
#include<cassert>
#include "../kernels.h"
class ActivationLeakyRT : public IPlugin {
public:
ActivationLeakyRT() {
}
~ActivationLeakyRT(){
}
int getNbOutputs() const override {
return 1;
}
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
return inputs[0];
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
size = 1;
for(int i=0; i<outputDims[0].nbDims; i++)
size *= outputDims[0].d[i];
}
int initialize() override {
return 0;
}
virtual void terminate() override {
}
virtual size_t getWorkspaceSize(int maxBatchSize) const override {
return 0;
}
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
activationLEAKYForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), size, stream);
return 0;
}
virtual size_t getSerializationSize() override {
return 1*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, size);
}
int size;
};
+168
View File
@@ -0,0 +1,168 @@
#include <vector>
#include <assert.h>
#include <algorithm>
#include <iterator>
#include "NvInfer.h"
class BatchStream
{
public:
BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches)
{
mBatchSize = batchSize;
mMaxBatches = maxBatches;
mDims = nvinfer1::DimsNCHW{ dim.n, dim.c, dim.h, dim.w };
mImageSize = mDims.c()*mDims.h()*mDims.w();
mBatch.resize(mBatchSize*mImageSize, 0);
mLabels.resize(mBatchSize, 0);
mFileBatch.resize(mDims.n()*mImageSize, 0);
mFileLabels.resize(mDims.n(), 0);
reset(0);
}
void reset(int firstBatch)
{
mBatchCount = 0;
mFileCount = 0;
mFileBatchPos = mDims.n();
skip(firstBatch);
}
bool next()
{
std::cout<<"Next batch: "<<mBatchCount<<" of "<<mMaxBatches<<"\n";
if (mBatchCount == mMaxBatches)
return false;
for (int csize = 1, batchPos = 0; batchPos < mBatchSize; batchPos += csize, mFileBatchPos += csize)
{
assert(mFileBatchPos > 0 && mFileBatchPos <= mDims.n());
if (mFileBatchPos == mDims.n() && !update())
return false;
// copy the smaller of: elements left to fulfill the request, or elements left in the file buffer.
csize = std::min(mBatchSize - batchPos, mDims.n() - mFileBatchPos);
std::copy_n(getFileBatch() + mFileBatchPos * mImageSize, csize * mImageSize, getBatch() + batchPos * mImageSize);
std::copy_n(getFileLabels() + mFileBatchPos, csize, getLabels() + batchPos);
}
mBatchCount++;
return true;
}
void skip(int skipCount)
{
if (mBatchSize >= mDims.n() && mBatchSize%mDims.n() == 0 && mFileBatchPos == mDims.n())
{
mFileCount += skipCount * mBatchSize / mDims.n();
std::cout<<mFileCount<<"\n";
return;
}
int x = mBatchCount;
for (int i = 0; i < skipCount; i++)
next();
mBatchCount = x;
}
float *getBatch() { return &mBatch[0]; }
float *getLabels() { return &mLabels[0]; }
int getBatchesRead() const { return mBatchCount; }
int getBatchSize() const { return mBatchSize; }
nvinfer1::DimsNCHW getDims() const { return mDims; }
private:
float* getFileBatch() { return &mFileBatch[0]; }
float* getFileLabels() { return &mFileLabels[0]; }
bool update()
{
std::string inputFileName = std::string("calibBatches/batch") + std::to_string(mFileCount++);
FILE * file = fopen(inputFileName.c_str(), "rb");
if (!file) {
FatalError("cant open batch calib file: " + inputFileName);
return false;
}
size_t readInputCount = fread(getFileBatch(), sizeof(float), mDims.n()*mImageSize, file);
size_t readLabelCount = fread(getFileLabels(), sizeof(float), mDims.n(), file);;
assert(readInputCount == size_t(mDims.n()*mImageSize) && readLabelCount == size_t(mDims.n()));
fclose(file);
mFileBatchPos = 0;
return true;
}
int mBatchSize{ 0 };
int mMaxBatches{ 0 };
int mBatchCount{ 0 };
int mFileCount{ 0 }, mFileBatchPos{ 0 };
int mImageSize{ 0 };
nvinfer1::DimsNCHW mDims;
std::vector<float> mBatch;
std::vector<float> mLabels;
std::vector<float> mFileBatch;
std::vector<float> mFileLabels;
};
class Int8EntropyCalibrator : public IInt8EntropyCalibrator
{
public:
Int8EntropyCalibrator(BatchStream& stream, int firstBatch, bool readCache = true)
: mStream(stream), mReadCache(readCache)
{
DimsNCHW dims = mStream.getDims();
mInputCount = mStream.getBatchSize() * dims.c() * dims.h() * dims.w();
checkCuda(cudaMalloc(&mDeviceInput, mInputCount * sizeof(float)));
mStream.reset(firstBatch);
}
virtual ~Int8EntropyCalibrator()
{
checkCuda(cudaFree(mDeviceInput));
}
int getBatchSize() const override { return mStream.getBatchSize(); }
bool getBatch(void* bindings[], const char* names[], int nbBindings) override
{
std::cout<<"CALIB request batch\n";
if (!mStream.next())
return false;
checkCuda(cudaMemcpy(mDeviceInput, mStream.getBatch(), mInputCount * sizeof(float), cudaMemcpyHostToDevice));
bindings[0] = mDeviceInput;
return true;
}
const void* readCalibrationCache(size_t& length) override
{
mCalibrationCache.clear();
std::ifstream input("table.calib", std::ios::binary);
input >> std::noskipws;
FatalError("rewrite different");
//if (mReadCache && input.good())
// std::copy(std::istream_iterator<char>(input), std::istream_iterator<char>(), std::back_inserter(mCalibrationCache));
length = mCalibrationCache.size();
return length ? &mCalibrationCache[0] : nullptr;
}
void writeCalibrationCache(const void* cache, size_t length) override
{
std::ofstream output("table.calib", std::ios::binary);
output.write(reinterpret_cast<const char*>(cache), length);
}
private:
BatchStream mStream;
bool mReadCache{ true };
size_t mInputCount;
void* mDeviceInput{ nullptr };
std::vector<char> mCalibrationCache;
};
+94
View File
@@ -0,0 +1,94 @@
#include<cassert>
#include "../kernels.h"
class RegionRT : public IPlugin {
public:
RegionRT(int classes, int coords, int num) {
this->classes = classes;
this->coords = coords;
this->num = num;
}
~RegionRT(){
}
int getNbOutputs() const override {
return 1;
}
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
return inputs[0];
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
c = inputDims[0].d[0];
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
int initialize() override {
return 0;
}
virtual void terminate() override {
}
virtual size_t getWorkspaceSize(int maxBatchSize) const override {
return 0;
}
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
for (int b = 0; b < batchSize; ++b){
for(int n = 0; n < num; ++n){
int index = entry_index(b, n*w*h, 0, batchSize);
activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream);
index = entry_index(b, n*w*h, coords, batchSize);
activationLOGISTICForward(srcData + index, dstData + index, w*h, stream);
}
}
//softmax start
int index = entry_index(0, 0, coords + 1, batchSize);
softmaxForward( srcData + index, classes, batchSize*num,
(batchSize*c*h*w)/num,
w*h, 1, w*h, 1, dstData + index, stream);
return 0;
}
virtual size_t getSerializationSize() override {
return 6*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, classes);
tk::dnn::writeBUF(buf, coords);
tk::dnn::writeBUF(buf, num);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
}
int c, h, w;
int classes, coords, num;
int entry_index(int batch, int location, int entry, int batchSize) {
int n = location / (w*h);
int loc = location % (w*h);
return batch*c*h*w*batchSize + n*w*h*(coords+classes+1) + entry*w*h + loc;
}
};
+63
View File
@@ -0,0 +1,63 @@
#include<cassert>
#include "../kernels.h"
class ReorgRT : public IPlugin {
public:
ReorgRT(int stride) {
this->stride = stride;
}
~ReorgRT(){
}
int getNbOutputs() const override {
return 1;
}
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
return DimsCHW{inputs[0].d[0]*stride*stride, inputs[0].d[1]/stride, inputs[0].d[2]/stride};
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
c = inputDims[0].d[0];
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
int initialize() override {
return 0;
}
virtual void terminate() override {
}
virtual size_t getWorkspaceSize(int maxBatchSize) const override {
return 0;
}
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
reorgForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]),
batchSize, c, h, w, stride, stream);
return 0;
}
virtual size_t getSerializationSize() override {
return 4*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, stride);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
}
int c, h, w, stride;
};
+82
View File
@@ -0,0 +1,82 @@
#include<cassert>
#include "../kernels.h"
class RouteRT : public IPlugin {
public:
RouteRT() {
}
~RouteRT(){
}
int getNbOutputs() const override {
return 1;
}
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
int out_c = 0;
for(int i=0; i<nbInputDims; i++) out_c += inputs[i].d[0];
return DimsCHW{out_c, inputs[0].d[1], inputs[0].d[2]};
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
in = nbInputs;
c = 0;
for(int i=0; i<nbInputs; i++) {
c_in[i] = inputDims[i].d[0];
c += inputDims[i].d[0];
}
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
int initialize() override {
return 0;
}
virtual void terminate() override {
}
virtual size_t getWorkspaceSize(int maxBatchSize) const override {
return 0;
}
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
int offset = 0;
for(int i=0; i<in; i++) {
dnnType *input = (dnnType*)reinterpret_cast<const dnnType*>(inputs[i]);
int in_dim = c_in[i]*h*w;
checkCuda( cudaMemcpyAsync(dstData + offset, input, in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) );
offset += in_dim;
}
return 0;
}
virtual size_t getSerializationSize() override {
return (4+MAX_INPUTS)*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, in);
for(int i=0; i<MAX_INPUTS; i++)
tk::dnn::writeBUF(buf, c_in[i]);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
}
static const int MAX_INPUTS = 4;
int in;
int c_in[MAX_INPUTS];
int c, h, w;
};
+65
View File
@@ -0,0 +1,65 @@
#include<cassert>
#include "../kernels.h"
class ShortcutRT : public IPlugin {
public:
ShortcutRT() {
}
~ShortcutRT(){
}
int getNbOutputs() const override {
return 1;
}
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
return DimsCHW{inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]};
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
c = inputDims[0].d[0];
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
int initialize() override {
return 0;
}
virtual void terminate() override {
}
virtual size_t getWorkspaceSize(int maxBatchSize) const override {
return 0;
}
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *srcDataBack = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, c, h, w, 1, stream);
return 0;
}
virtual size_t getSerializationSize() override {
return 3*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
}
int c, h, w;
};
+65
View File
@@ -0,0 +1,65 @@
#include<cassert>
#include "../kernels.h"
class UpsampleRT : public IPlugin {
public:
UpsampleRT(int stride) {
this->stride = stride;
}
~UpsampleRT(){
}
int getNbOutputs() const override {
return 1;
}
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
return DimsCHW(inputs[0].d[0], inputs[0].d[1]*stride, inputs[0].d[2]*stride);
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
c = inputDims[0].d[0];
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
int initialize() override {
return 0;
}
virtual void terminate() override {
}
virtual size_t getWorkspaceSize(int maxBatchSize) const override {
return 0;
}
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
fill(dstData, batchSize*c*h*w*stride*stride, 0.0, stream);
upsampleForward(srcData, dstData, batchSize, c, h, w, stride, 1, 1, stream);
return 0;
}
virtual size_t getSerializationSize() override {
return 4*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, stride);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
}
int c, h, w, stride;
};
+116
View File
@@ -0,0 +1,116 @@
#include<cassert>
#include "../kernels.h"
#define YOLORT_CLASSNAME_W 256
class YoloRT : public IPlugin {
public:
YoloRT(int classes, int num, tk::dnn::Yolo *yolo = nullptr) {
this->classes = classes;
this->num = num;
mask = new dnnType[num];
bias = new dnnType[num*3*2];
if(yolo != nullptr) {
memcpy(mask, yolo->mask_h, sizeof(dnnType)*num);
memcpy(bias, yolo->bias_h, sizeof(dnnType)*num*3*2);
classesNames = yolo->classesNames;
}
}
~YoloRT(){
}
int getNbOutputs() const override {
return 1;
}
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
return inputs[0];
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
c = inputDims[0].d[0];
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
int initialize() override {
return 0;
}
virtual void terminate() override {
}
virtual size_t getWorkspaceSize(int maxBatchSize) const override {
return 0;
}
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
for (int b = 0; b < batchSize; ++b){
for(int n = 0; n < num; ++n){
int index = entry_index(b, n*w*h, 0, batchSize);
activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream);
index = entry_index(b, n*w*h, 4, batchSize);
activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream);
}
}
//std::cout<<"YOLO END\n";
return 0;
}
virtual size_t getSerializationSize() override {
return 5*sizeof(int) + num*sizeof(dnnType) + num*3*2*sizeof(dnnType) + YOLORT_CLASSNAME_W*classes*sizeof(char);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, classes);
tk::dnn::writeBUF(buf, num);
tk::dnn::writeBUF(buf, c);
tk::dnn::writeBUF(buf, h);
tk::dnn::writeBUF(buf, w);
for(int i=0; i<num; i++)
tk::dnn::writeBUF(buf, mask[i]);
for(int i=0; i<3*2*num; i++)
tk::dnn::writeBUF(buf, bias[i]);
// save classes names
for(int i=0; i<classes; i++) {
char tmp[YOLORT_CLASSNAME_W];
strcpy(tmp, classesNames[i].c_str());
for(int j=0; j<YOLORT_CLASSNAME_W; j++) {
tk::dnn::writeBUF(buf, tmp[j]);
}
}
}
int c, h, w;
int classes, num;
std::vector<std::string> classesNames;
dnnType *mask;
dnnType *bias;
int entry_index(int batch, int location, int entry, int batchSize) {
int n = location / (w*h);
int loc = location % (w*h);
return batch*c*h*w*batchSize + n*w*h*(4+classes+1) + entry*w*h + loc;
}
};
+2 -10
View File
@@ -3,14 +3,6 @@
*/
#include "Network.h"
#include "Layer.h"
#include "NetworkRT.h"
namespace tkDNN {
/**
Return the tkDNN version
*/
int getVersion() {
return 100;
}
}
#define TKDNN_VERSION 400
+116
View File
@@ -0,0 +1,116 @@
#ifndef UTILS_H
#define UTILS_H
#include <iostream>
#include <sstream>
#include <fstream>
#include <iomanip>
#include <stdlib.h>
#include "cuda.h"
#include "cuda_runtime_api.h"
#include <cublas_v2.h>
#include <cudnn.h>
#define dnnType float
// Colored output
#define COL_END "\033[0m"
#define COL_RED "\033[31m"
#define COL_GREEN "\033[32m"
#define COL_ORANGE "\033[33m"
#define COL_BLUE "\033[34m"
#define COL_PURPLE "\033[35m"
#define COL_CYAN "\033[36m"
#define COL_REDB "\033[1;31m"
#define COL_GREENB "\033[1;32m"
#define COL_ORANGEB "\033[1;33m"
#define COL_BLUEB "\033[1;34m"
#define COL_PURPLEB "\033[1;35m"
#define COL_CYANB "\033[1;36m"
// Simple Timer
#define TIMER_START \
timespec start, end; \
clock_gettime(CLOCK_MONOTONIC, &start);
#define TIMER_STOP_C(col) \
clock_gettime(CLOCK_MONOTONIC, &end); \
double t_ns = ((double)(end.tv_sec - start.tv_sec) * 1.0e9 + \
(double)(end.tv_nsec - start.tv_nsec)) / \
1.0e6; \
std::cout << col << "Time:" << std::setw(16) << t_ns << " ms\n" \
<< COL_END;
#define TIMER_STOP TIMER_STOP_C(COL_CYANB)
/********************************************************
* Prints the error message, and exits
* ******************************************************/
#define EXIT_WAIVED 0
#define FatalError(s) \
{ \
std::stringstream _where, _message; \
_where << __FILE__ << ':' << __LINE__; \
_message << std::string(s) + "\n" \
<< __FILE__ << ':' << __LINE__; \
std::cerr << _message.str() << "\nAborting...\n"; \
cudaDeviceReset(); \
exit(EXIT_FAILURE); \
}
#define checkCUDNN(status) \
{ \
std::stringstream _error; \
if (status != CUDNN_STATUS_SUCCESS) \
{ \
_error << "CUDNN failure: " << cudnnGetErrorString(status); \
FatalError(_error.str()); \
} \
}
#define checkCuda(status) \
{ \
std::stringstream _error; \
if (status != 0) \
{ \
_error << "Cuda failure: " << cudaGetErrorString(status); \
FatalError(_error.str()); \
} \
}
#define checkERROR(status) \
{ \
std::stringstream _error; \
if (status != 0) \
{ \
_error << "Generic failure: " << status; \
FatalError(_error.str()); \
} \
}
#define checkNULL(ptr) \
{ \
std::stringstream _error; \
if (ptr == nullptr) \
{ \
_error << "Null pointer"; \
FatalError(_error.str()); \
} \
}
void printCenteredTitle(const char *title, char fill, int dim);
bool fileExist(const char *fname);
void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek = 0);
int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device = true, int limit = 10);
void printDeviceVector(int size, dnnType* vec_d, bool device = true);
void resize(int size, dnnType **data);
void matrixTranspose(cublasHandle_t handle, dnnType *srcData, dnnType *dstData, int rows, int cols);
void matrixMulAdd(cublasHandle_t handle, dnnType *srcData, dnnType *dstData,
dnnType *add_vector, int dim, dnnType mul);
#endif //UTILS_H

Some files were not shown because too many files have changed in this diff Show More