diff --git a/.gitignore b/.gitignore
index c1d362c..22f7f94 100644
--- a/.gitignore
+++ b/.gitignore
@@ -19,4 +19,6 @@ demo/BDD100K_val
cmake-build-minsizerel/*
scripts/COCO_val2017/*
scripts/COCO_val2017.zip
-scripts/all_labels.txt
\ No newline at end of file
+scripts/all_labels.txt
+/cmake/cuda_script
+/cmake-build-debug/
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 877b506..c29e987 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -1,15 +1,69 @@
cmake_minimum_required(VERSION 3.15)
-
-project (tkDNN)
+project(tkDNN)
set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake)
-if(UNIX)
-set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++14 -fPIC -Wno-deprecated-declarations")
-endif()
-if(WIN32)
set(CMAKE_CXX_STANDARD 14)
-set(CMAKE_CXX_FLAGS "/O2 /FS /EHsc")
+
+option(ENABLE_OPENCV_CUDA_CONTRIB "Enable OpenCV CUDA Contrib" OFF )
+
+if(NOT CMAKE_BUILD_TYPE)
+ set(CMAKE_BUILD_TYPE "Release" CACHE STRING "default build" FORCE)
+endif(NOT CMAKE_BUILD_TYPE)
+
+find_package(CUDA 9.0 REQUIRED)
+if (CUDA_FOUND)
+ set(OUTPUTFILE ${CMAKE_CURRENT_SOURCE_DIR}/cmake/cuda_script) # No suffix required
+ execute_process(COMMAND "rm ${OUTPUTFILE}")
+ set(CUDAFILE ${CMAKE_CURRENT_SOURCE_DIR}/cmake/getCudaArch.cu)
+ execute_process(COMMAND ${CUDA_NVCC_EXECUTABLE} -lcuda ${CUDAFILE} -o ${OUTPUTFILE})
+ execute_process(COMMAND ${OUTPUTFILE}
+ RESULT_VARIABLE CUDA_RETURN_CODE
+ OUTPUT_VARIABLE ARCH)
+
+ if(${CUDA_RETURN_CODE} EQUAL 0)
+ set(CUDA_SUCCESS "TRUE")
+ else()
+ set(CUDA_SUCCESS "FALSE")
+ endif()
+
+ if (${CUDA_SUCCESS})
+ message(STATUS "CUDA Architecture: ${ARCH}")
+ message(STATUS "CUDA Version: ${CUDA_VERSION_STRING}")
+ message(STATUS "CUDA Path: ${CUDA_TOOLKIT_ROOT_DIR}")
+ message(STATUS "CUDA Libararies: ${CUDA_LIBRARIES}")
+ message(STATUS "CUDA Performance Primitives: ${CUDA_npp_LIBRARY}")
+ set(CUDA_NVCC_FLAGS "${ARCH}")
+ else()
+ message(WARNING ${ARCH})
+ endif()
+endif()
+
+SET(CUDA_SEPARABLE_COMPILATION ON)
+
+if(UNIX)
+ if(CMAKE_BUILD_TYPE MATCHES Release)
+ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fPIC -Wno-deprecated-declarations -Wno-unused-variable -O3")
+ set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} --maxrregcount=32)
+ endif()
+
+ if(CMAKE_BUILD_TYPE MATCHES Debug)
+ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fPIC -Wno-deprecated-declarations -Wno-unused-variable -g3")
+ set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} --maxrregcount=32 -G -g)
+ endif()
+endif()
+
+if(WIN32)
+ if(CMAKE_BUILD_TYPE MATCHES Release)
+ set(CMAKE_CXX_FLAGS "/O2 /FS /EHsc /MD")
+ set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} --maxrregcount=32)
+ endif()
+
+ if(CMAKE_BUILD_TYPE MATCHES Debug)
+ set(CMAKE_CXX_FLAGS "/Od /FS /EHsc /MDd")
+ set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} --maxrregcount=32 -G -g)
+ endif()
set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON)
endif(WIN32)
+
include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN)
# project specific flags
@@ -18,7 +72,7 @@ if(DEBUG)
endif()
if(TKDNN_PATH)
- message("SET TKDNN_PATH:"${TKDNN_PATH})
+ message("SET TKDNN_PATH:" ${TKDNN_PATH})
add_definitions(-DTKDNN_PATH="${TKDNN_PATH}")
else()
add_definitions(-DTKDNN_PATH="${CMAKE_CURRENT_SOURCE_DIR}")
@@ -28,20 +82,21 @@ endif()
#-------------------------------------------------------------------------------
# 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'")
-set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} --maxrregcount=32)
+set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS}" --compiler-options '-fPIC')
+
find_package(CUDNN REQUIRED)
include_directories(${CUDNN_INCLUDE_DIR})
+find_package(yaml-cpp REQUIRED)
+
# compile
-file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/sorting.cu")
+file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/sorting.cu" "src/pluginsRT/*.cpp")
cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIRS})
cuda_add_library(kernels SHARED ${tkdnn_CUSRC})
-target_link_libraries(kernels ${CUDA_CUBLAS_LIBRARIES})
+target_link_libraries(kernels ${CUDA_CUBLAS_LIBRARIES} ${CUDA_LIBRARIES} ${CUDNN_LIBRARIES} yaml-cpp)
+
#-------------------------------------------------------------------------------
@@ -53,12 +108,23 @@ include_directories(${EIGEN3_INCLUDE_DIR})
find_package(OpenCV REQUIRED)
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV")
+if(ENABLE_OPENCV_CUDA_CONTRIB)
+ if (OpenCV_FOUND)
+ find_package(OpenCV COMPONENTS cudawarping cudaarithm)
+ if(OpenCV_cudawarping_FOUND AND OpenCV_cudaarithm_FOUND)
+ add_compile_definitions(OPENCV_CUDACONTRIB)
+ message("OpenCV Cuda Contrib modules found")
+ else()
+ message("OpenCV Cuda Contrib modules not found")
+ set(ENABLE_OPENCV_CUDA_CONTRIB OFF)
+ endif()
+ endif()
+endif()
# if(OpenCV_CUDA_VERSION)
# add_compile_definitions(OPENCV_CUDACONTRIB)
# endif()
# gives problems in cross-compiling, probably malformed cmake config
-find_package(yaml-cpp REQUIRED)
#-------------------------------------------------------------------------------
# Build Libraries
@@ -69,7 +135,7 @@ set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDNN_LIBRAR
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS} ${NVINFER_INCLUDES})
add_library(tkDNN SHARED ${tkdnn_SRC})
-target_link_libraries(tkDNN ${tkdnn_LIBS})
+target_link_libraries(tkDNN ${tkdnn_LIBS} ${CUDA_CUBLAS_LIBRARIES})
#static
#add_library(tkDNN_static STATIC ${tkdnn_SRC})
@@ -143,6 +209,14 @@ target_link_libraries(test_shelfnet_mapillary tkDNN)
add_executable(test_shelfnet_coco tests/shelfnet/shelfnet_coco.cpp)
target_link_libraries(test_shelfnet_coco tkDNN)
+# MONODEPTH2
+add_executable(test_monodepth2_640 tests/monodepth2/monodepth2_640.cpp)
+target_link_libraries(test_monodepth2_640 tkDNN)
+
+add_executable(test_monodepth2_1024 tests/monodepth2/monodepth2_1024.cpp)
+target_link_libraries(test_monodepth2_1024 tkDNN)
+
+
# DEMOS
add_executable(test_rtinference tests/test_rtinference/rtinference.cpp)
target_link_libraries(test_rtinference tkDNN)
@@ -162,6 +236,9 @@ target_link_libraries(demoTracker tkDNN)
add_executable(seg_demo demo/demo/seg_demo.cpp)
target_link_libraries(seg_demo tkDNN)
+add_executable(demoDepth demo/demo/demoDepth.cpp)
+target_link_libraries(demoDepth tkDNN)
+
#-------------------------------------------------------------------------------
# Install
#-------------------------------------------------------------------------------
@@ -171,7 +248,7 @@ target_link_libraries(seg_demo tkDNN)
#endif()
message("install dir:" ${CMAKE_INSTALL_PREFIX})
install(DIRECTORY include/ DESTINATION include/)
-install(TARGETS tkDNN kernels DESTINATION lib)
+install(TARGETS tkDNN DESTINATION lib)
install(TARGETS test_simple test_mnist test_mnistRT test_rtinference demo map_demo DESTINATION bin)
install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/cmake/" # source directory
DESTINATION "share/tkDNN/cmake/" # target directory
diff --git a/README.md b/README.md
index ddfd4ef..825c00b 100644
--- a/README.md
+++ b/README.md
@@ -23,9 +23,11 @@ If you use tkDNN in your research, please cite the [following paper](https://iee
- [x] Support 2D/3D Object Detection and Tracking [README](docs/README_2d3dtracking.md)
#### 24 November 2021
- [x] Support to sematic segmentation on cuda 11
-- [x] Support to TensorRT8 (thanks to [Harshvardhan Chandirasekar](https://github.com/perseusdg)).
+- [x] Support to TensorRT8. (thanks to [Harshvardhan Chandirasekar](https://github.com/perseusdg))
+#### 30 March 2022
+- [x] Support to monocular depth esitmation [README](docs/README_depth.md) (thanks to [Harshvardhan Chandirasekar](https://github.com/perseusdg))
+
-TensorRT8 (and therefore Jetpack 4.6) is currently supported only on the branch tensorrt8 due to [performance issue with TensorRT8](https://docs.nvidia.com/deeplearning/tensorrt/release-notes/tensorrt-8.html)). We will merge it to the master as soon as those issues are fixed (probably in future minor releases).
## FPS Results
Inference FPS of yolov4 with tkDNN, average of 1200 images with the same dimension as the input size, on
@@ -80,17 +82,17 @@ Results for COCO val 2017 (5k images), on RTX 2080Ti, with conf threshold=0.001
- [Workflow](#workflow)
- [Exporting weights](#exporting-weights)
- [Run the demos](#run-the-demos)
- - [tkDNN on Windows 10 (experimental)](#tkdnn-on-windows-10-experimental)
+ - [tkDNN on Windows 10 or Windows 11](#tkdnn-on-windows-10-or-windows-11)
- [Existing tests and supported networks](#existing-tests-and-supported-networks)
- [References](#references)
## Dependencies
This branch works on every NVIDIA GPU that supports the following (latest tested) dependencies:
-* CUDA 11.0 (or >= 10) [the segmentation only works with CUDA 10 for now]
-* cuDNN 8.0.4 (or >= 7.3)
-* TensorRT 7.2.0 (or >=5)
-* OpenCV 4.5.2 (or >=4)
+* CUDA 11.3 (or >= 10.2)
+* cuDNN 8.2.1 (or >= 8.0.4)
+* TensorRT 8.0.3 (or >=7.2)
+* OpenCV 4.5.4 (or >=4)
* cmake 3.21 (or >= 3.15)
* yaml-cpp 0.5.2
* eigen3 3.3.4
@@ -106,16 +108,18 @@ To compile and install OpenCV4 with contrib us the script ```install_OpenCV4.sh`
```
bash scripts/install_OpenCV4.sh
```
-When using openCV not compiled with contrib, comment the definition of OPENCV_CUDACONTRIBCONTRIB in include/tkDNN/DetectionNN.h. When commented, the preprocessing of the networks is computed on the CPU, otherwise on the GPU. In the latter case some milliseconds are saved in the end-to-end latency.
+If you have OpenCV compiled with cuda and contrib and want to use it with tkDNN pass ```ENABLE_OPENCV_CUDA_CONTRIB=ON``` flag when compiling tkDBB
+. If the flag is not passed,the preprocessing of the networks is computed on the CPU, otherwise on the GPU. In the latter case some milliseconds are saved in the end-to-end latency.
## How to compile this repo
-Build with cmake. If using Ubuntu 18.04 a new version of cmake is needed (3.15 or above).
+Build with cmake. If using Ubuntu 18.04 a new version of cmake is needed (3.15 or above).
+On both linux and windows ,the ```CMAKE_BUILD_TYPE``` variable needs to be defined as either ```Release``` or ```Debug```.
```
git clone https://github.com/ceccocats/tkDNN
cd tkDNN
mkdir build
cd build
-cmake ..
+cmake -DCMAKE_BUILD_TYPE=Release ..
make
```
@@ -136,14 +140,15 @@ For specific details on how to export weights see [HERE](./docs/exporting_weight
For specific details on how to run:
- 2D object detection demos, details on FP16, INT8 and batching see [HERE](./docs/demo.md).
- segmentation demos see [HERE](./docs/README_seg.md).
+- monocular depth estimation see [HERE](./docs/README_depth.md).
- 2D/3D object detection and tracking demos see [HERE](./docs/README_2d3dtracking.md).
- mAP demo to evaluate 2D object detectors see [HERE](./docs/mAP_demo.md).

-## tkDNN on Windows 10 (experimental)
+## tkDNN on Windows 10 or Windows 11
-For specific details on how to run tkDNN on Windows 10 see [HERE](./docs/windows.md).
+For specific details on how to run tkDNN on Windows 10/11 see [HERE](./docs/windows.md).
## Existing tests and supported networks
@@ -182,6 +187,8 @@ For specific details on how to run tkDNN on Windows 10 see [HERE](./docs/windows
| shelfnet_berkeley | ShelfNet18_realtime11 | [DeepDrive](https://bdd-data.berkeley.edu/) | 20 | 1024x1024 | [weights](https://cloud.hipert.unimore.it/s/m92e7QdD9gYMF7f/download) |
| dla34_cnet3d | Centernet3D (DLA34 backend)4 | [KITTI 2017](http://www.cvlibs.net/datasets/kitti/eval_object.php?obj_benchmark=3d) | 1 | 512x512 | [weights](https://cloud.hipert.unimore.it/s/2MDyWGzQsTKMjmR/download) |
| dla34_ctrack | CenterTrack (DLA34 backend)12 | [NuScenes 3D](https://www.nuscenes.org/) | 7 | 512x512 | [weights](https://cloud.hipert.unimore.it/s/rjNfgGL9FtAXLHp/download) |
+| monodepth2 | Monodepth2 13 | [KITTI DEPTH](http://www.cvlibs.net/datasets/kitti/raw_data.php) | - | 640x192 | [weights-mono](https://cloud.hipert.unimore.it/s/iYw9QwgP6CsqxLR/download) |
+| monodepth2 | Monodepth2 13 | [KITTI DEPTH](http://www.cvlibs.net/datasets/kitti/raw_data.php) | - | 640x192 | [weights-stereo](https://cloud.hipert.unimore.it/s/XmwbWNXDfqyQ4EL/download) |
## References
@@ -198,3 +205,11 @@ For specific details on how to run tkDNN on Windows 10 see [HERE](./docs/windows
10. Wang, Chien-Yao, Alexey Bochkovskiy, and Hong-Yuan Mark Liao. "Scaled-YOLOv4: Scaling Cross Stage Partial Network." arXiv preprint arXiv:2011.08036 (2020).
11. Zhuang, Juntang, et al. "ShelfNet for fast semantic segmentation." Proceedings of the IEEE International Conference on Computer Vision Workshops. 2019.
12. Zhou, Xingyi, Vladlen Koltun, and Philipp Krähenbühl. "Tracking objects as points." European Conference on Computer Vision. Springer, Cham, 2020.
+13. Godard, Clément, et al. "Digging into self-supervised monocular depth estimation." Proceedings of the IEEE/CVF International Conference on Computer Vision. 2019.
+
+## Contributors
+The main contibutors, in chronological order, are:
+- [Francesco Gatti](https://github.com/ceccocats), francesco.gatti@hipert.it
+- [Micaela Verucchi](https://github.com/mive93), micaela.verucchi@unimore.it
+- [Davide Sapienza](https://github.com/sapienzadavide), davide.sapienza@unimore.it
+- [Harshvardhan Chandirasekar](https://github.com/perseusdg), f20180523@goa.bits-pilani.ac.in
diff --git a/cmake/getCudaArch.cu b/cmake/getCudaArch.cu
new file mode 100644
index 0000000..1d66199
--- /dev/null
+++ b/cmake/getCudaArch.cu
@@ -0,0 +1,20 @@
+#include
+
+int main(int argc, char **argv){
+ cudaDeviceProp dP;
+ float min_cc = 5.0;
+
+ int rc = cudaGetDeviceProperties(&dP, 0);
+ if(rc != cudaSuccess) {
+ cudaError_t error = cudaGetLastError();
+ printf("CUDA error: %s", cudaGetErrorString(error));
+ return rc; /* Failure */
+ }
+ if((dP.major+(dP.minor/10)) < min_cc) {
+ printf("Min Compute Capability of %2.1f required: %d.%d found\n Not Building CUDA Code", min_cc, dP.major, dP.minor);
+ return 1; /* Failure */
+ } else {
+ printf("-arch=sm_%d%d", dP.major, dP.minor);
+ return 0; /* Success */
+ }
+}
\ No newline at end of file
diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp
index 5445d86..c857086 100644
--- a/demo/demo/demo.cpp
+++ b/demo/demo/demo.cpp
@@ -37,14 +37,18 @@ int main(int argc, char *argv[]) {
if(!fileExist(net.c_str()))
FatalError("The given network does not exist. Create the rt first.");
- #ifdef __linux__
+ #ifdef __linux__
std::string input = YAMLgetConf(conf, "input", "../demo/yolo_test.mp4");
+ std::string cfgPath = YAMLgetConf(conf,"cfg_input", "../tests/darknet/cfg/yolo4tiny.cfg");
+ std::string namePath = YAMLgetConf(conf,"name_input","../tests/darknet/names/coco.names");
#elif _WIN32
- std::string input = YAMLgetConf(conf, "win_input", "..\\..\\..\\demo\\yolo_test.mp4");
+ std::string input = YAMLgetConf(conf, "win_input", "..\\..\\..\\demo\\yolo_test.mp4");
+ std::string cfgPath = YAMLgetConf(conf,"cfg_win_input","..\\..\\..\\tests\\darknet\\cfg\\yolo4tiny.cfg");
+ std::string namePath = YAMLgetConf(conf,"name_win_input","..\\..\\..\\tests\\darknet\\names\\coco.names");
#endif
- if(!fileExist(input.c_str()))
+ if(!fileExist(input.c_str()))
FatalError("The given input video does not exist.");
-
+
char ntype = YAMLgetConf(conf, "ntype", 'y');
int n_classes = YAMLgetConf(conf, "n_classes", 80);
int n_batch = YAMLgetConf(conf, "n_batch", 1);
@@ -66,7 +70,7 @@ int main(int argc, char *argv[]) {
// create detection network
tk::dnn::Yolo3Detection yolo;
tk::dnn::CenternetDetection cnet;
- tk::dnn::MobilenetDetection mbnet;
+ tk::dnn::MobilenetDetection mbnet;
tk::dnn::DetectionNN *detNN;
@@ -86,7 +90,12 @@ int main(int argc, char *argv[]) {
FatalError("Network type not allowed (3rd parameter)\n");
}
- detNN->init(net, n_classes, n_batch, conf_thresh);
+ if(ntype == 'c' || ntype == 'm'){
+ cfgPath = "";
+ namePath = "";
+ }
+
+ detNN->init(net,cfgPath,namePath,n_classes,n_batch,conf_thresh);
// open video stream
cv::VideoCapture cap(input);
@@ -146,10 +155,10 @@ int main(int argc, char *argv[]) {
double mean = 0;
std::cout<stats.begin(), detNN->stats.end())/n_batch<<" ms\n";
- std::cout<<"Max: "<<*std::max_element(detNN->stats.begin(), detNN->stats.end())/n_batch<<" ms\n";
+ std::cout<<"Min: "<<*std::min_element(detNN->stats.begin(), detNN->stats.end())<<" ms\n";
+ std::cout<<"Max: "<<*std::max_element(detNN->stats.begin(), detNN->stats.end())<<" ms\n";
for(int i=0; istats.size(); i++) mean += detNN->stats[i]; mean /= detNN->stats.size();
- std::cout<<"Avg: "<
+#include
+#include /* srand, rand */
+//#include
+#include
+
+#include "tkDNN/DepthNN.h"
+
+bool gRun;
+
+void sig_handler(int signo) {
+ std::cout<<"request gateway stop\n";
+ gRun = false;
+}
+
+int main(int argc, char *argv[]) {
+
+ signal(SIGINT, sig_handler);
+
+ std::string net = "monodepth2_fp32.rt";
+ if(argc > 1)
+ net = argv[1];
+ #ifdef __linux__
+ std::string input = "../demo/yolo_test.mp4";
+ #elif _WIN32
+ std::string input = "..\\..\\..\\demo\\yolo_test.mp4";
+ #endif
+ if(argc > 2)
+ input = argv[2];
+ bool show = true;
+ if(argc > 3)
+ show = atoi(argv[3]);
+ bool save = true;
+ if(argc > 4)
+ save = atoi(argv[4]);
+
+ std::cout <<"Net settings - net: "<< net
+ <<"\n";
+ std::cout <<"Demo settings - input: "<< input
+ <<", show: "<< show
+ <<", save: "<< save<<"\n\n";
+
+ tk::dnn::DepthNN depthNN;
+
+ // create depth network
+ int n_batch = 1;
+ depthNN.init(net, n_batch);
+
+ // open video stream
+ cv::VideoCapture cap(input);
+ if(!cap.isOpened())
+ gRun = false;
+ else
+ std::cout<<"camera started\n";
+
+ cv::VideoWriter resultVideo;
+ if(save) {
+ int w = depthNN.output_w;
+ int h = depthNN.output_h;
+ resultVideo.open("result.mp4", cv::VideoWriter::fourcc('M','P','4','V'), 30, cv::Size(w, h));
+ }
+
+ if(show)
+ cv::namedWindow("depth", cv::WINDOW_NORMAL);
+
+ cv::Mat frame;
+ std::vector batch_frame;
+ std::vector batch_dnn_input;
+
+ // start detection loop
+ gRun = true;
+ while(gRun) {
+ batch_dnn_input.clear();
+ batch_frame.clear();
+
+ //read frame
+ cap >> frame;
+ if(!frame.data)
+ break;
+ batch_frame.push_back(frame);
+ batch_dnn_input.push_back(frame.clone());
+
+ //inference
+ depthNN.update(batch_dnn_input, 1);
+ if(show){
+ cv::imshow("depth", depthNN.depthMats[0]);
+ cv::waitKey(1);
+
+ }
+
+ if(save)
+ resultVideo << depthNN.depthMats[0];
+ }
+
+ std::cout<<"detection end\n";
+
+ double mean = 0;
+ std::cout< 1)
net = argv[1];
if(argc > 2)
- ntype = argv[2][0];
+ ntype = argv[2][0];
if(argc > 3)
- labels_path = argv[3];
+ cfg_path = argv[3];
if(argc > 4)
- config_filename = argv[4];
+ name_path = argv[4];
if(argc > 5)
- n_batches = atoi(argv[5]);
+ labels_path = argv[5];
if(argc > 6)
- confidence_thresh = atof(argv[6]);
+ config_filename = argv[6];
+ if(argc > 7)
+ n_batches = atoi(argv[7]);
+ if(argc > 8)
+ confidence_thresh = atof(argv[8]);
std::cout<<"conf t: "<init(net, n_classes, 1, conf_thresh);
+ detNN->init(net,cfg_path,name_path,n_classes, 1, conf_thresh);
//read images
std::ifstream all_labels(labels_path);
diff --git a/demo/demo/seg_demo.cpp b/demo/demo/seg_demo.cpp
index 2b491d8..f22e5c2 100644
--- a/demo/demo/seg_demo.cpp
+++ b/demo/demo/seg_demo.cpp
@@ -1,7 +1,9 @@
#include
#include
#include /* srand, rand */
+#ifdef __linux__
#include
+#endif
#include
#include "SegmentationNN.h"
diff --git a/demo/demoConfig.yaml b/demo/demoConfig.yaml
index 6cdaff7..194d466 100644
--- a/demo/demoConfig.yaml
+++ b/demo/demoConfig.yaml
@@ -2,6 +2,14 @@
input : "../demo/yolo_test.mp4"
win_input : "..\\..\\..\\demo\\yolo_test.mp4"
+#cfg input
+cfg_input : "../tests/darknet/cfg/yolo4tiny.cfg"
+cfg_win_input : "..\\..\\..\\tests\\darknet\\cfg\\yolo4tiny.cfg"
+
+#name input
+name_input : "../tests/darknet/names/coco.names"
+name_win_input : "..\\..\\..\\tests\\darknet\\names\\coco.names"
+
# network config
net : "yolo4tiny_fp32.rt"
ntype : 'y'
@@ -11,4 +19,4 @@ conf_thresh : 0.3
# demo config
show : true
-save : true
\ No newline at end of file
+save : false
diff --git a/docker/Dockerfile.base b/docker/Dockerfile.base
index e61b0d3..1e235d1 100644
--- a/docker/Dockerfile.base
+++ b/docker/Dockerfile.base
@@ -1,57 +1,140 @@
-FROM nvidia/cuda:10.2-cudnn7-devel-ubuntu18.04
-LABEL maintainer "Francesco Gatti"
+FROM nvidia/cudagl:11.3.1-devel-ubuntu20.04
-ADD nv-tensorrt-repo-ubuntu1804-cuda10.2-trt7.0.0.11-ga-20191216_1-1_amd64.deb /tmp/trt.deb
-RUN apt-get update && dpkg -i /tmp/trt.deb && rm /tmp/trt.deb && apt-get update
-RUN apt install -y libnvinfer7=7.0.0-1+cuda10.2 libnvinfer-dev=7.0.0-1+cuda10.2
-RUN DEBIAN_FRONTEND=noninteractive apt install -y git wget libeigen3-dev libyaml-cpp-dev
-RUN cd /tmp && \
- wget https://github.com/Kitware/CMake/releases/download/v3.17.3/cmake-3.17.3-Linux-x86_64.sh && \
- chmod +x cmake-3.17.3-Linux-x86_64.sh && \
- ./cmake-3.17.3-Linux-x86_64.sh --prefix=/usr/local --exclude-subdir --skip-license && \
- rm ./cmake-3.17.3-Linux-x86_64.sh
+LABEL maintainer "TKDNN AUTHORS"
+LABEL Description="tkDNN+cudagl"
+LABEL com.tkdnn.nvidia.version="11.3.1"
-RUN echo "INSTALL OPENCV"
-RUN apt-get install -y build-essential \
- unzip \
- pkg-config \
- libjpeg-dev \
- libpng-dev \
- libtiff-dev \
- libavcodec-dev \
- libavformat-dev \
- libswscale-dev \
- libv4l-dev \
- libxvidcore-dev \
- libx264-dev \
- libgtk-3-dev \
- libatlas-base-dev \
- gfortran \
- libgstreamer1.0-dev \
- libgstreamer-plugins-base1.0-dev \
- libdc1394-22-dev \
- libavresample-dev
-RUN cd && wget https://github.com/opencv/opencv/archive/4.3.0.tar.gz && tar -xf 4.3.0.tar.gz && rm *.tar.gz
-RUN cd && wget https://github.com/opencv/opencv_contrib/archive/4.3.0.tar.gz && tar -xf 4.3.0.tar.gz && rm *.tar.gz
-RUN cd && \
- cd opencv-4.3.0 && mkdir build && cd build && \
+ENV DEBIAN_FRONTEND noninteractive
+ENV CC gcc
+ENV CXX g++
+
+RUN apt-get update && apt-get install -y \
+libblkid-dev && apt-get clean && rm -rf /var/lib/apt/lists/*
+
+RUN apt-get update && apt-get install -y \
+libcudnn8-dev=8.2.1.32-1+cuda11.3 \
+libcudnn8=8.2.1.32-1+cuda11.3 \
+libnvinfer-dev=8.0.3-1+cuda11.3 \
+libnvinfer8=8.0.3-1+cuda11.3 && apt-get clean && rm -rf /var/lib/apt/lists/*
+
+RUN apt-get update && apt-get install -y --no-install-recommends \
+libblkid-dev \
+locales \
+lsb-release \
+mesa-utils \
+git \
+nano \
+terminator \
+wget \
+curl \
+libssl-dev \
+htop \
+dbus-x11 \
+libqt5opengl5-dev \
+libgtk-3-dev \
+libvtk7-dev \
+libv4l-dev \
+tar \
+libgoogle-glog-dev \
+libgflags-dev \
+gfortran-9 \
+libtbb-dev \
+libgstreamer1.0-dev \
+libgstreamer-plugins-base1.0-dev \
+libdc1394-22-dev \
+libavresample-dev \
+libatlas-cpp-0.6-dev \
+python3-dev \
+gdb \
+python3-pip \
+unzip libtbb-dev && \
+apt-get clean && rm -rf /var/lib/apt/lists/*
+
+RUN apt-get update && apt-get install -y --no-install-recommends \
+software-properties-common && apt-get clean && rm -rf /var/lib/apt/lists/*
+
+RUN apt-add-repository universe
+RUN apt-get update && apt-get install -y python3-pip python3 openssh-server ssh pyqt5-dev sip-dev && apt-get clean && rm -rf /var/lib/apt/lists/*
+RUN pip3 install --upgrade pip
+RUN pip3 install --upgrade virtualenv
+RUN pip3 install --upgrade paramiko
+RUN pip3 install --ignore-installed --upgrade numpy protobuf
+
+
+RUN cd ~ && mkdir build
+RUN cd ~/build && wget https://github.com/Kitware/CMake/releases/download/v3.21.4/cmake-3.21.4.tar.gz && \
+tar -xvf cmake-3.21.4.tar.gz && cd cmake-3.21.4 && ./configure --prefix=/usr/local --qt-gui --parallel=12 && \
+make -j8 && make install
+
+RUN apt-get update && apt-get install -y automake autoconf pkg-config libevent-dev libncurses5-dev bison && \
+apt-get clean && rm -rf /var/lib/apt/lists/
+
+RUN git clone https://github.com/tmux/tmux.git && \
+cd tmux && git checkout tags/3.2 && ls -la && sh autogen.sh && ./configure && make -j8 && make install
+
+RUN apt-get update && apt-get install -y zsh && apt-get clean && rm -rf /var/lib/apt/lists/*
+RUN wget https://github.com/robbyrussell/oh-my-zsh/raw/master/tools/install.sh -O - | zsh || true
+RUN chsh -s /usr/bin/zsh root
+RUN git clone https://github.com/sindresorhus/pure /root/.oh-my-zsh/custom/pure
+RUN ln -s /root/.oh-my-zsh/custom/pure/pure.zsh-theme /root/.oh-my-zsh/custom/
+RUN ln -s /root/.oh-my-zsh/custom/pure/async.zsh /root/.oh-my-zsh/custom/
+RUN sed -i -e 's/robbyrussell/refined/g' /root/.zshrc
+RUN sed -i '/plugins=(/c\plugins=(git git-flow adb pyenv tmux)' /root/.zshrc
+
+RUN mkdir -p /root/.config/terminator/
+COPY assets/terminator_config /root/.config/terminator/config
+
+RUN echo "/usr/local/nvidia/lib" >> /etc/ld.so.conf.d/nvidia.conf && \
+ echo "/usr/local/nvidia/lib64" >> /etc/ld.so.conf.d/nvidia.conf && \
+ echo "/usr/local/cuda/lib64" >> /etc/ld.so.conf.d/nvidia.conf
+
+
+ENV PATH /usr/local/nvidia/bin:/usr/local/cuda/bin:${PATH}
+ENV LD_LIBRARY_PATH /usr/local/nvidia/lib:/usr/local/nvidia/lib64:/usr/local/cuda/lib64:/usr/lib:/usr/lib/x86_64-linux-gnu:/usr/local/lib:${LD_LIBRARY_PATH}
+ENV NVIDIA_VISIBLE_DEVICES all
+ENV NVIDIA_DRIVER_CAPABILITIES compute,utility,graphics
+
+
+
+RUN cd ~/build && wget https://github.com/opencv/opencv/archive/4.5.4.tar.gz && tar -xf 4.5.4.tar.gz && rm 4.5.4.tar.gz
+RUN cd ~/build && wget https://github.com/opencv/opencv_contrib/archive/4.5.4.tar.gz && tar -xf 4.5.4.tar.gz && rm 4.5.4.tar.gz
+RUN cd ~/build && \
+ cd opencv-4.5.4 && mkdir build && cd build && \
cmake -D CMAKE_BUILD_TYPE=RELEASE \
-D CMAKE_INSTALL_PREFIX=/usr/local \
-D INSTALL_PYTHON_EXAMPLES=OFF \
-D INSTALL_C_EXAMPLES=OFF \
- -D OPENCV_EXTRA_MODULES_PATH='~/opencv_contrib-4.3.0/modules' \
+ -D OPENCV_EXTRA_MODULES_PATH='~/build/opencv_contrib-4.5.4/modules' \
-D BUILD_EXAMPLES=OFF \
+ -D BUILD_TESTS=OFF \
+ -D BUILD_PERF_TESTS=OFF \
+ -D BUILD_DOCS=OFF \
-D WITH_CUDA=ON \
+ -D WITH_OPENGL=ON \
+ -D WITH_NVCUVID=ON \
-D CUDA_ARCH_BIN=7.2 \
- -D CUDA_ARCH_PTX="" \
+ -D CUDA_ARCH_PTX=7.2 \
-D ENABLE_FAST_MATH=ON \
-D CUDA_FAST_MATH=ON \
-D WITH_CUBLAS=ON \
+ -D WITH_CUDNN=ON \
+ -D WITH_OPENMP=ON \
+ -D WITH_NONFREE=ON \
-D WITH_LIBV4L=ON \
-D WITH_GSTREAMER=ON \
-D WITH_GSTREAMER_0_10=OFF \
-D WITH_TBB=ON \
- ../ && make -j12 && make install
-RUN apt clean
+ ../ && make -j12 && make install && ldconfig
+RUN cd ~ && rm -rf build
+RUN cd ~ && mkdir Development && cd Development && \
+git clone https://github.com/ceccocats/tkDNN.git && cd tkDNN && \
+mkdir build && cd build && \
+cmake -DCMAKE_BUILD_TYPE=Release .. && \
+make -j6
+
+RUN apt-get clean && rm -rf /var/lib/apt/lists/*
+COPY assets/entrypoint_setup.sh /
+ENTRYPOINT ["/entrypoint_setup.sh"]
+CMD ["terminator"]
\ No newline at end of file
diff --git a/docker/README.md b/docker/README.md
index aec202a..15f3987 100644
--- a/docker/README.md
+++ b/docker/README.md
@@ -9,13 +9,10 @@ docker build -t tkdnn:build -f Dockerfile .
# make nvidia docker working
# follow this guide: https://github.com/NVIDIA/nvidia-docker
-# dowload tensorrt
-# from: https://developer.nvidia.com/compute/machine-learning/tensorrt/secure/7.0/7.0.0.11/local_repo/nv-tensorrt-repo-ubuntu1804-cuda10.2-trt7.0.0.11-ga-20191216_1-1_amd64.deb
-
# build image
docker build -t ceccocats/tkdnn:latest -f Dockerfile.base .
# run image
-docker run -ti --gpus all --rm ceccocats/tkdnn:latest bash
+./docker_launch.sh
```
diff --git a/docker/assets/entrypoint_setup.sh b/docker/assets/entrypoint_setup.sh
new file mode 100755
index 0000000..348a4f5
--- /dev/null
+++ b/docker/assets/entrypoint_setup.sh
@@ -0,0 +1,123 @@
+#! /bin/bash
+
+CMD=
+
+# Functions
+# TOOD: Check if we can use: getent passwd $USER to extract all variables
+# TODO: Check for valid inputs, cause now it will go through even with bad inputs
+check_envs () {
+ DOCKER_CUSTOM_USER_OK=true;
+ if [ -z ${DOCKER_USER_NAME+x} ]; then
+ DOCKER_CUSTOM_USER_OK=false;
+ return;
+ fi
+
+ if [ -z ${DOCKER_USER_ID+x} ]; then
+ DOCKER_CUSTOM_USER_OK=false;
+ return;
+ else
+ if ! [ -z "${DOCKER_USER_ID##[0-9]*}" ]; then
+ echo -e "\033[1;33mWarning: User-ID should be a number. Falling back to defaults.\033[0m"
+ DOCKER_CUSTOM_USER_OK=false;
+ return;
+ fi
+ fi
+
+ if [ -z ${DOCKER_USER_GROUP_NAME+x} ]; then
+ DOCKER_CUSTOM_USER_OK=false;
+ return;
+ fi
+
+ if [ -z ${DOCKER_USER_GROUP_ID+x} ]; then
+ DOCKER_CUSTOM_USER_OK=false;
+ return;
+ else
+ if ! [ -z "${DOCKER_USER_GROUP_ID##[0-9]*}" ]; then
+ echo -e "\033[1;33mWarning: Group-ID should be a number. Falling back to defaults.\033[0m"
+ DOCKER_CUSTOM_USER_OK=false;
+ return;
+ fi
+ fi
+}
+
+setup_env_user () {
+ USER=$1
+ USER_ID=$2
+ GROUP=$3
+ GROUP_ID=$4
+
+ ## Create user
+ useradd -m $USER
+
+ ## Copy zsh/sh configs
+ cp /root/.profile /home/$USER/
+ cp /root/.bashrc /home/$USER/
+ cp /root/.zshrc /home/$USER/
+ ## Copy terminator configs
+ mkdir -p /home/$USER/.config/terminator
+ cp /root/.config/terminator/config /home/$USER/.config/terminator/config
+ cp /root/.config/terminator/background.png /home/$USER/.config/terminator/background.png
+ cp -rf /root/.oh-my-zsh /home/$USER/
+ cp -rf /root/tkDNN /home/$USER/
+ rm -rf /home/$USER/.oh-my-zsh/custom/pure.zsh-theme /home/$USER/.oh-my-zsh/custom/async.zsh
+ ln -s /home/$USER/.oh-my-zsh/custom/pure/pure.zsh-theme /home/$USER/.oh-my-zsh/custom/
+ ln -s /home/$USER/.oh-my-zsh/custom/pure/async.zsh /home/$USER/.oh-my-zsh/custom/
+ sed -i -e 's@ZSH=\"/root@ZSH=\"/home/$USER@g' /home/$USER/.zshrc
+ # Copy SSH keys & fix owner
+ if [ -d "/root/.ssh" ]; then
+ cp -rf /root/.ssh /home/$USER/
+ chown -R $USER:$GROUP /home/$USER/.ssh
+ fi
+
+ ## Fix owner
+ chown $USER:$GROUP /home/$USER
+ chown -R $USER:$GROUP /home/$USER/.config
+ chown $USER:$GROUP /home/$USER/.profile
+ chown $USER:$GROUP /home/$USER/.bashrc
+ chown $USER:$GROUP /home/$USER/.zshrc
+ chown -R $USER:$GROUP /home/$USER/.oh-my-zsh
+ chown -R $USER:$GROUP /home/$USER/tkDNN
+
+ ## This a trick to keep the evnironmental variables of root which is important!
+ echo "if ! [ \"$DOCKER_USER_NAME\" = \"$(id -un)\" ]; then" >> /root/.bashrc
+ echo " cd /home/$DOCKER_USER_NAME" >> /root/.bashrc
+ echo " su $DOCKER_USER_NAME" >> /root/.bashrc
+ echo "fi" >> /root/.bashrc
+
+ echo "if ! [ \"$DOCKER_USER_NAME\" = \"$(id -un)\" ]; then" >> /root/.zshrc
+ echo " cd /home/$DOCKER_USER_NAME" >> /root/.zshrc
+ echo " su $DOCKER_USER_NAME" >> /root/.zshrc
+ echo "fi" >> /root/.zshrc
+
+ ## Setup Password-file
+ PASSWDCONTENTS=$(grep -v "^${USER}:" /etc/passwd)
+ GROUPCONTENTS=$(grep -v -e "^${GROUP}:" -e "^docker:" /etc/group)
+
+ (echo "${PASSWDCONTENTS}" && echo "${USER}:x:$USER_ID:$GROUP_ID::/home/$USER:/bin/bash") > /etc/passwd
+ (echo "${GROUPCONTENTS}" && echo "${GROUP}:x:${GROUP_ID}:") > /etc/group
+ (if test -f /etc/sudoers ; then echo "${USER} ALL=(ALL) NOPASSWD: ALL" >> /etc/sudoers ; fi)
+}
+
+
+# ---Main---
+
+# Create new user
+## Check Inputs
+check_envs
+
+## Determine user & Setup Environment
+if [ $DOCKER_CUSTOM_USER_OK == true ]; then
+ echo " -->DOCKER_USER Input is set to '$DOCKER_USER_NAME:$DOCKER_USER_ID:$DOCKER_USER_GROUP_NAME:$DOCKER_USER_GROUP_ID'";
+ echo -e "\033[0;32mSetting up environment for user=$DOCKER_USER_NAME\033[0m"
+ setup_env_user $DOCKER_USER_NAME $DOCKER_USER_ID $DOCKER_USER_GROUP_NAME $DOCKER_USER_GROUP_ID
+else
+ echo " -->DOCKER_USER* variables not set. Using 'root'.";
+ echo -e "\033[0;32mSetting up environment for user=root\033[0m"
+ DOCKER_USER_NAME="root"
+fi
+
+# Change shell to zsh
+chsh -s /usr/bin/zsh $DOCKER_USER_NAME
+
+# Run CMD from Docker
+"$@"
\ No newline at end of file
diff --git a/docker/assets/terminator_config b/docker/assets/terminator_config
new file mode 100644
index 0000000..d65a1b3
--- /dev/null
+++ b/docker/assets/terminator_config
@@ -0,0 +1,18 @@
+[global_config]
+ title_transmit_bg_color = "#2e3436"
+[keybindings]
+[layouts]
+ [[default]]
+ [[[child1]]]
+ parent = window0
+ type = Terminal
+ [[[window0]]]
+ parent = ""
+ type = Window
+[plugins]
+[profiles]
+ [[default]]
+ background_color = "#282828"
+ cursor_color = "#aaaaaa"
+ foreground_color = "#f3f3f3"
+ palette = "#000000:#aa0000:#00aa00:#c4a000:#3465a4:#75507b:#06989a:#d3d7cf:#88807c:#f15d22:#73c48f:#ffce51:#48b9c7:#ad7fa8:#34e2e2:#eeeeec"
diff --git a/docker/docker_launch.sh b/docker/docker_launch.sh
new file mode 100755
index 0000000..24adb53
--- /dev/null
+++ b/docker/docker_launch.sh
@@ -0,0 +1,9 @@
+xhost local:root
+docker run --rm -it --runtime=nvidia --privileged --net=host --cap-add sys_ptrace -d --ipc=host \
+-v /tmp/.X11-unix:/tmp/.X11-unix -e DISPLAY=$DISPLAY \
+-v $HOME/.Xauthority:/home/$(id -un)/.Xauthority -e XAUTHORITY=/home/$(id -un)/.Xauthority \
+-e DOCKER_USER_NAME=$(id -un) \
+-e DOCKER_USER_ID=$(id -u) \
+-e DOCKER_USER_GROUP_NAME=$(id -gn) \
+-e DOCKER_USER_GROUP_ID=$(id -g) \
+-v $HOME/.ssh:/home/$(id -un)/.ssh ceccocats/tkdnn
diff --git a/docs/README_depth.md b/docs/README_depth.md
new file mode 100644
index 0000000..ac87bb6
--- /dev/null
+++ b/docs/README_depth.md
@@ -0,0 +1,54 @@
+# Monocular depth estimation with tkDNN
+
+Currently tkDNN supports only Monodepth2 as monocular depth esitmation network.
+
+
+## Run the demo
+
+To run the depth estimation demo follow these steps (example with monodepth2):
+```
+rm monodepth2_fp32.rt # be sure to delete(or move) old tensorRT files
+./test_monodepth2 # run the yolo test (is slow)
+./demoDepth monodepth2_fp32.rt ../demo/yolo_test.mp4
+```
+In general the demo program takes the following parameters:
+```
+./demoDepth
+```
+where
+* `````` is the rt file generated by a test
+* ```<``` is the path to a video file or a camera input
+* `````` if set to 0 the demo will not show the visualization, it will otherwise (default=1)
+* `````` if set to 1 the demo will save the video into result.mp4, it won't otherwise (default=1)
+
+NB) By default it is used FP32 inference
+
+
+
+
+
+
+
diff --git a/docs/demo.md b/docs/demo.md
index 5971781..bf6b79d 100644
--- a/docs/demo.md
+++ b/docs/demo.md
@@ -26,26 +26,28 @@ rm yolo4_fp32.rt # be sure to delete(or move) old tensorRT files
```
If you get problems in the creation, try to check the error activating the debug of TensorRT in this way:
```
-cmake .. -DDEBUG=True
+cmake .. -DCMAKE_BUILD_TYPE=Debug -DDEBUG=True
make
```
-Once you have successfully created your rt file, run the demo:
+Once you have successfully created your rt file, run the demo:
```
-./demo
+./ demo
```
-In general the demo program takes 1 parameter, the `````` that is the path to che configuration file. The parameter is optional and its default value is ```"../demo/demoConfig.yaml"```.
+In general the demo program takes 1 parameter, the `````` that is the path to che configuration file. The parameter is optional and its default value is ```"../demo/demoConfig.yaml"```.
The config file is a yaml file with the following attributes:
-* ```net``` is the rt file generated by a test
-* ```input``` is the path to a video file or a camera input (on Linux)
-* ```win_input``` is the path to a video file or a camera input (on Windows)
-* ```ntype``` is the type of network. Thee types are currently supported: ```y``` (YOLO family), ```c``` (CenterNet family) and ```m``` (MobileNet-SSD family)
-* ```n_classes``` is the number of classes the network is trained on
-* ```n_batch``` number of batches to use in inference (N.B. you should first export TKDNN_BATCHSIZE to the required n_batches and create again the rt file for the network).
-* ```conf_thresh``` confidence threshold for the detector. Only bounding boxes with threshold greater than conf-thresh will be displayed.
-* ```show``` if set to 0 the demo will not show the visualization (if n-batches ==1)
-* ```save``` if set to 1 the demo will save the video of the demo into result.mp4 (if n-batches ==1)
+* ```net``` is the rt file generated by a test
+* ```input``` is the path to a video file or a camera input (on Linux)
+* ```win_input``` is the path to a video file or a camera input (on Windows)
+* ```ntype``` is the type of network. Thee types are currently supported: ```y``` (YOLO family), ```c``` (CenterNet family) and ```m``` (MobileNet-SSD family)
+* ```n_classes``` is the number of classes the network is trained on
+* ```n_batch``` number of batches to use in inference (N.B. you should first export TKDNN_BATCHSIZE to the required n_batches and create again the rt file for the network).
+* ```conf_thresh``` confidence threshold for the detector. Only bounding boxes with threshold greater than conf-thresh will be displayed.
+* ```show``` if set to 0 the demo will not show the visualization (if n-batches ==1)
+* ```save``` if set to 1 the demo will save the video of the demo into result.mp4 (if n-batches ==1)
+* ```cfg_input``` (for linux) \ ```cfg_win_input``` (for windows) is the location of the cfg path of the network for mobilenet and centernet networks use ```" "```
+* ```name_input``` (for linux) \ ```name_win_input``` (for windows) is the location of the name path of the network for mobilenet and centernet networks use ```" "```
N.B. By default it is used FP32 inference
@@ -58,10 +60,10 @@ N.B. By default it is used FP32 inference
To run the demo with FP16 inference follow these steps (example with yolov3):
```
export TKDNN_MODE=FP16 # set the half floating point optimization
-rm yolo3_fp16.rt # be sure to delete(or move) old tensorRT files
-./test_yolo3 # run the yolo test (is slow)
-# set net: yolo3_fp16.rt in the config-file
-./demo
+rm yolo4_fp16.rt # be sure to delete(or move) old tensorRT files
+./test_yolo4 # run the yolo test (is slow)
+#set net: yolo4_fp16.rt in the config file
+./demo
```
N.B. Using FP16 inference will lead to some errors in the results (first or second decimal).
@@ -84,10 +86,10 @@ Then a complete example using yolo3 and COCO dataset would be:
export TKDNN_MODE=INT8
export TKDNN_CALIB_LABEL_PATH=../demo/COCO_val2017/all_labels.txt
export TKDNN_CALIB_IMG_PATH=../demo/COCO_val2017/all_images.txt
-rm yolo3_int8.rt # be sure to delete(or move) old tensorRT files
-./test_yolo3 # run the yolo test (is slow)
-# set net: yolo3_int8.rt in the config-file
-./demo
+rm yolo4_int8.rt # be sure to delete(or move) old tensorRT files
+./test_yolo4 # run the yolo test (is slow)
+#set net: yolo4_int8.rt in the config file
+./demo
```
N.B.
diff --git a/docs/exporting_weights.md b/docs/exporting_weights.md
index 811431d..b4cb366 100644
--- a/docs/exporting_weights.md
+++ b/docs/exporting_weights.md
@@ -86,6 +86,18 @@ mkdir layer debug
python export.py
```
+### 6)Export weights for monodepth2
+To get the weights needed to run Shelfnet tests use [this](https://github.com/perseusdg/monodepth2) fork of a Pytorch implementation of monodepth2 network.
+
+```
+git clone https://github.com/perseusdg/monodepth2
+cd monodepth2
+mkdir models # Download the official weights and put depth.pth and encorder.pth inside this new folder
+conda env create --file monodepth.yaml
+conda activate monodepth2
+python exporter.py # you will find the weights inside the tkDNN_bin folder
+```
+
## Darknet Parser
tkDNN implement and easy parser for darknet cfg files, a network can be converted with *tk::dnn::darknetParser*:
```
diff --git a/docs/windows.md b/docs/windows.md
index 60813c5..fced442 100644
--- a/docs/windows.md
+++ b/docs/windows.md
@@ -7,17 +7,18 @@
- [Run the demo on Windows](#run-the-demo-on-windows)
- [FP16 inference windows](#fp16-inference-windows)
- [INT8 inference windows](#int8-inference-windows)
+ - [Run tkDNN on WSL2 with cuda](#tkdnn-on-cuda-wsl)
- [Known issues with tkDNN on Windows](#known-issues-with-tkdnn-on-windows)
### Dependencies-Windows
This branch should work on every NVIDIA GPU supported in windows with the following dependencies:
-* WINDOWS 10 1803 or HIGHER
-* CUDA 10.0 (Recommended CUDA 11.2 )
-* CUDNN 7.6 (Recommended CUDNN 8.1.1 )
-* TENSORRT 6.0.1 (Recommended TENSORRT 7.2.3.4 )
-* OPENCV 3.4 (Recommended OPENCV 4.2.0 )
-* MSVC 16.7
+* WINDOWS 10 1803/WINDOWS 11 or HIGHER
+* CUDA 11.2
+* CUDNN 8.1.1
+* TENSORRT 7.2.3
+* OPENCV 4.2
+* MSVC 16.9+
* YAML-CPP
* EIGEN3
* 7ZIP (ADD TO PATH)
@@ -58,7 +59,7 @@ To run the object detection file create .rt file bu running:
Once the rt file has been successfully create,run the demo using the following command:
```
-.\demo.exe yolo4tiny_fp32.rt ..\demo\yolo_test.mp4 y
+.\demo.exe yolo4_fp32.rt ..\demo\yolo_test.mp4 y 80 ..\tests\darknet\cfg\yolo4.cfg ..\tests\darknet\names\cococ.names
```
For general info on more demo paramters,check Run the demo section on top
To run the test_all_tests.sh on windows,use git bash or msys2
@@ -85,11 +86,17 @@ del /f yolo4tiny_int8.rt # be sure to delete(or move) old tensorRT files
```
+### Run tkDNN on WSL2 with cuda
+tkDNN works on wsl2 with cuda,although not all networks (centernet,mobilenet) work properly.
+If you encounter issues with running the network as a result of driver not found or cuda launch error,running the following command should solve the issue
+```cp /usr/lib/wsl/lib/lib* /usr/lib/x86_64-linux-gnu/ ```
+
+
+
### Known issues with tkDNN on Windows
-Mobilenet and Centernet demos work properly only when built with msvc 16.7 in Release Mode,when built in debug mode for the mentioned networks one might encounter opencv assert errors
+In theory all models (centernet,mobilenet,darknet,centertrack,cnet3d and shelfnet) should work on Windows.
-All Darknet models work properly with demo using MSVC version(16.7-16.9)
+On pascal cards(sm 6x) ,nvidia cuda wsl driver 510.06 don't work well with tkDNN both on windows and cuda wsl , Nvidia drivers >465+ and < 500 are completely supported .
-It is recommended to use Nvidia Driver(465+),Cuda unknown errors have been observed when using older drivers on pascal(SM 61) devices.
diff --git a/include/tkDNN/CenterTrack.h b/include/tkDNN/CenterTrack.h
index 783a44e..aa573fa 100644
--- a/include/tkDNN/CenterTrack.h
+++ b/include/tkDNN/CenterTrack.h
@@ -13,6 +13,11 @@
#include "TrackingNN.h"
+#ifdef _WIN32
+#define _USE_MATH_DEFINES
+#include
+#endif
+
#include "kernelsThrust.h"
diff --git a/include/tkDNN/CenternetDetection.h b/include/tkDNN/CenternetDetection.h
index 3c8cfbb..07c80cd 100644
--- a/include/tkDNN/CenternetDetection.h
+++ b/include/tkDNN/CenternetDetection.h
@@ -73,7 +73,7 @@ public:
CenternetDetection() {};
~CenternetDetection() {};
- bool init(const std::string& tensor_path, const int n_classes=80, const int n_batches=1, const float conf_thresh=0.3);
+ bool init(const std::string& tensor_path,const std::string& cfg_path,const std::string& name_path, const int n_classes=80, const int n_batches=1, const float conf_thresh=0.3);
void preprocess(cv::Mat &frame, const int bi=0);
void postprocess(const int bi=0,const bool mAP=false);
};
diff --git a/include/tkDNN/CenternetDetection3D.h b/include/tkDNN/CenternetDetection3D.h
index f9c918f..9f2b214 100644
--- a/include/tkDNN/CenternetDetection3D.h
+++ b/include/tkDNN/CenternetDetection3D.h
@@ -9,6 +9,11 @@
#include // std::iota
#include // std::sort
+#ifdef _WIN32
+#define _USE_MATH_DEFINES
+#include
+#endif
+
#include "DetectionNN3D.h"
#include "kernelsThrust.h"
diff --git a/include/tkDNN/DarknetParser.h b/include/tkDNN/DarknetParser.h
index 089c4d6..c6d2472 100644
--- a/include/tkDNN/DarknetParser.h
+++ b/include/tkDNN/DarknetParser.h
@@ -47,5 +47,8 @@ namespace tk { namespace dnn {
std::vector &netLayers, const std::vector& names);
std::vector darknetReadNames(const std::string& names_file);
tk::dnn::Network* darknetParser(const std::string& cfg_file, const std::string& wgs_path, const std::string& names_file);
+ void loadYoloInfo(const std::string &cfg_file,int lineNo,std::vector &mask,std::vector &anchors,int &num,int &classes,float &nms_thresh,int &nms_kind,int &coords);
+ void loadYoloInitInfo(int &channels,int &width,int &height,const std::string &cfg_file);
+ std::vector noYolosLine(const std::string &cfg_file);
}}
diff --git a/include/tkDNN/DepthNN.h b/include/tkDNN/DepthNN.h
new file mode 100644
index 0000000..4d34239
--- /dev/null
+++ b/include/tkDNN/DepthNN.h
@@ -0,0 +1,180 @@
+#ifndef DEPTHNN_H
+#define DEPTHNN_H
+
+#include
+#include
+#include
+#ifdef __linux__
+#include
+#endif
+
+#include
+
+#include
+#include
+#include
+
+#include "tkDNN/utils.h"
+#include "tkDNN/tkdnn.h"
+
+#include "NetworkViz.h"
+
+
+namespace tk { namespace dnn {
+
+class DepthNN {
+
+ public:
+ tk::dnn::NetworkRT *netRT = nullptr;
+ dnnType *input_h;
+ dnnType *input_d;
+ float* depth_h;
+
+ int output_w;
+ int output_h;
+
+ int nBatches = 1;
+
+ cv::Mat bgr[3];
+ cv::Mat imagePreproc;
+
+ std::vector stats; /*keeps track of inference times (ms)*/
+ std::vector> depths;
+ std::vector depthMats;
+
+ DepthNN() {};
+ ~DepthNN(){};
+
+ /**
+ * Method used to initialize the class, allocate memory and compute
+ * needed data.
+ *
+ * @param tensor_path path to the rt file of the NN.
+ * @param n_batches maximum number of batches to use in inference
+ * @return true if everything is correct, false otherwise.
+ */
+ void init(const std::string& tensor_path, const int n_batches=1){
+ //create net
+
+ std::cout<<(tensor_path).c_str()<<"\n";
+ nBatches = n_batches;
+ netRT = new tk::dnn::NetworkRT(NULL, (tensor_path).c_str());
+
+ //allocate memory for NN input
+ checkCuda(cudaMallocHost(&input_h, sizeof(dnnType) * netRT->input_dim.tot() * nBatches));
+ checkCuda(cudaMalloc(&input_d, sizeof(dnnType) * netRT->input_dim.tot() * nBatches));
+
+ //allocate memory for NN output
+ depthMats.resize(nBatches);
+ depths.resize(nBatches);
+ for(int i=0; i< depths.size();++i)
+ depths[i].resize(netRT->buffersDIM[1].tot());
+
+ depth_h = (float *)malloc(netRT->buffersDIM[1].tot() * sizeof(float));
+
+ output_h = netRT->buffersDIM[1].h;
+ output_w = netRT->buffersDIM[1].w;
+
+ }
+
+
+ /**
+ * This method preprocess the image, before feeding it to the NN.
+ *
+ * @param frame original frame to adapt for inference.
+ * @param bi batch index
+ */
+ void preprocess(cv::Mat &frame, const int bi=0) {
+ //resize image, remove mean, divide by std
+ cv::Mat frame_nomean;
+ resize(frame, frame, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
+ frame.convertTo(frame_nomean, CV_32FC3);
+ frame_nomean.convertTo(imagePreproc, CV_32FC3, 1 / 255.0, 0);
+
+ //copy image into tensor and copy it into GPU
+ cv::split(imagePreproc, bgr);
+ for (int i = 0; i < netRT->input_dim.c; i++){
+ int idx = i * imagePreproc.rows * imagePreproc.cols;
+ int ch = netRT->input_dim.c-1 -i;
+ memcpy((void *)&input_h[idx + netRT->input_dim.tot()*bi], (void *)bgr[ch].data, imagePreproc.rows * imagePreproc.cols * sizeof(dnnType));
+ }
+ checkCuda(cudaMemcpyAsync(input_d+ netRT->input_dim.tot()*bi, input_h + netRT->input_dim.tot()*bi, netRT->input_dim.tot() * sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
+ }
+
+ /**
+ * This method postprocess the output of the NN to obtain the correct
+ * boundig boxes.
+ *
+ * @param bi batch index
+ * @param mAP set to true only if all the probabilities for a bounding
+ * box are needed, as in some cases for the mAP calculation
+ */
+ void postprocess(const int bi=0) {
+
+ dnnType *rt_out[1];
+ rt_out[0] = (dnnType *)netRT->buffersRT[1]+ netRT->buffersDIM[1].tot()*bi;
+ checkCuda(cudaMemcpy(depth_h, rt_out[0], netRT->buffersDIM[1].tot()* sizeof(float), cudaMemcpyDeviceToHost));
+ memcpy(&depths[bi][0], &depth_h[0], netRT->buffersDIM[1].tot()* sizeof(float));
+
+ // cv::Mat d(netRT->buffersDIM[1].h, netRT->buffersDIM[1].w, CV_8UC1, depth_h);
+ // depthMats[bi] = d.clone();
+
+ cv::Mat depth_mat = vizData2Mat(rt_out[0], netRT->buffersDIM[1], netRT->buffersDIM[1].h, netRT->buffersDIM[1].w);
+ // cv::Mat depth_mat = vizData2Mat((dnnType *)netRT->buffersRT[0], netRT->buffersDIM[0], netRT->buffersDIM[0].h, netRT->buffersDIM[0].w);
+ depthMats[bi] = depth_mat.clone();
+
+ }
+
+ /**
+ * This method performs the inference of the NN.
+ *
+ * @param frames frames to build the embedding from.
+ * @param cur_batches number of batches to use in inference
+ */
+ void update(std::vector& frames, const int cur_batches=1){
+ if(cur_batches > nBatches)
+ FatalError("A batch size greater than nBatches cannot be used");
+
+ if(TKDNN_VERBOSE) printCenteredTitle(" TENSORRT feature extraction ", '=', 30);
+ {
+ TKDNN_TSTART
+ for(int bi=0; biinput_dim;
+ dim.n = cur_batches;
+ {
+ if(TKDNN_VERBOSE) dim.print();
+ TKDNN_TSTART
+ netRT->infer(dim, input_d);
+ TKDNN_TSTOP
+ if(TKDNN_VERBOSE) dim.print();
+ stats.push_back(t_ns);
+ }
+
+ {
+ TKDNN_TSTART
+ for(int bi=0; bi& mListIn);
@@ -55,7 +55,7 @@ private:
int mFileBatchPos{ 0 };
int mImageSize{ 0 };
- nvinfer1::DimsNCHW mDims;
+ nvinfer1::Dims4 mDims;
std::vector mBatch;
std::vector mLabels;
std::vector mFileBatch;
diff --git a/include/tkDNN/Int8Calibrator.h b/include/tkDNN/Int8Calibrator.h
index 4a0ea47..c7a7d97 100644
--- a/include/tkDNN/Int8Calibrator.h
+++ b/include/tkDNN/Int8Calibrator.h
@@ -30,10 +30,10 @@ public:
Int8EntropyCalibrator(BatchStream& stream, int firstBatch, const std::string& calibTableFilePath,
const std::string& inputBlobName, bool readCache = true);
virtual ~Int8EntropyCalibrator() { checkCuda(cudaFree(mDeviceInput)); }
- int getBatchSize() const override { return mStream.getBatchSize(); }
- bool getBatch(void* bindings[], const char* names[], int nbBindings) override;
- const void* readCalibrationCache(size_t& length) override;
- void writeCalibrationCache(const void* cache, size_t length) override;
+ int getBatchSize() const NOEXCEPT override { return mStream.getBatchSize(); }
+ bool getBatch(void* bindings[], const char* names[], int nbBindings) NOEXCEPT override;
+ const void* readCalibrationCache(size_t& length) NOEXCEPT override;
+ void writeCalibrationCache(const void* cache, size_t length) NOEXCEPT override;
private:
BatchStream mStream;
diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h
index d1234a5..daa27e5 100644
--- a/include/tkDNN/Layer.h
+++ b/include/tkDNN/Layer.h
@@ -31,7 +31,8 @@ enum layerType_t {
LAYER_SHORTCUT,
LAYER_UPSAMPLE,
LAYER_REGION,
- LAYER_YOLO
+ LAYER_YOLO,
+ LAYER_PADDING,
};
#define TKDNN_BN_MIN_EPSILON 1e-5
@@ -56,8 +57,8 @@ public:
int id = 0;
bool final; //if the layer is the final one
- uint n_params = 0;
- uint feature_map_size = 0;
+ unsigned int n_params = 0;
+ unsigned int feature_map_size = 0;
long unsigned MACC = 0;
@@ -87,6 +88,7 @@ public:
case LAYER_UPSAMPLE: return "Upsample";
case LAYER_REGION: return "Region";
case LAYER_YOLO: return "Yolo";
+ case LAYER_PADDING: return "Padding";
default: return "unknown";
}
}
@@ -423,6 +425,8 @@ public:
virtual layerType_t getLayerType() { return LAYER_FLATTEN; };
virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+
+ int c, h, w, rows, cols;
};
/**
@@ -436,6 +440,7 @@ public:
virtual layerType_t getLayerType() { return LAYER_RESHAPE; };
virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+ int n,c,h,w;
};
@@ -470,7 +475,6 @@ public:
virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
-protected:
dnnType mul, add;
dnnType *add_vector;
};
@@ -497,6 +501,7 @@ public:
int winH, winW;
int strideH, strideW;
int paddingH, paddingW;
+ int padding;
bool size;
tkdnnPoolingMode_t pool_mode;
@@ -516,9 +521,35 @@ protected:
bool poolOn3d;
};
+/**
+ * Padding Layers
+ * tkDNN supports reflection,constant and zero padding
+ */
+
+typedef enum {
+ PADDING_MODE_CONSTANT = 0,
+ PADDING_MODE_ZERO = 1,
+ PADDING_MODE_REFLECTION = 2
+} tkdnnPaddingMode_t;
+
+class Padding : public Layer {
+public:
+ Padding(Network *net,int32_t pad_h,int32_t pad_w,tkdnnPaddingMode_t padding_mode,float constant = 0.0);
+ virtual ~Padding();
+ virtual layerType_t getLayerType(){return LAYER_PADDING ;};
+ virtual dnnType* infer(dataDim_t& dim,dnnType* srcData);
+ int32_t paddingH,paddingW;
+ tkdnnPaddingMode_t padding_mode;
+ float constant;
+
+};
+
+
+
/**
Softmax layer
*/
+
class Softmax : public Layer {
public:
@@ -582,6 +613,8 @@ public:
virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+ int c,h,w;
+
public:
Layer *backLayer;
bool mul = false;
@@ -602,6 +635,7 @@ public:
int stride;
bool reverse;
+ int c,h,w;
};
struct box {
@@ -685,6 +719,7 @@ public:
virtual layerType_t getLayerType() { return LAYER_REGION; };
int classes, coords, num;
+ int c,h,w;
virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
};
diff --git a/include/tkDNN/MobilenetDetection.h b/include/tkDNN/MobilenetDetection.h
index 9a5fedc..ec35b20 100644
--- a/include/tkDNN/MobilenetDetection.h
+++ b/include/tkDNN/MobilenetDetection.h
@@ -65,7 +65,7 @@ public:
MobilenetDetection() {};
~MobilenetDetection() {};
- bool init(const std::string& tensor_path, const int n_classes, const int n_batches=1, const float conf_thresh=0.3);
+ bool init(const std::string& tensor_path, const std::string& cfg_path,const std::string& name_path,const int n_classes, const int n_batches=1, const float conf_thresh=0.3);
void preprocess(cv::Mat &frame, const int bi=0);
void postprocess(const int bi=0,const bool mAP=false);
};
diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h
index 9892a24..a7e67c2 100644
--- a/include/tkDNN/NetworkRT.h
+++ b/include/tkDNN/NetworkRT.h
@@ -7,47 +7,30 @@
#include "Layer.h"
#include "NvInfer.h"
#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+#include
+
+
namespace tk { namespace dnn {
-template void writeBUF(char*& buffer, const T& val)
-{
- *reinterpret_cast(buffer) = val;
- buffer += sizeof(T);
-}
-
-template T readBUF(const char*& buffer)
-{
- T val = *reinterpret_cast(buffer);
- buffer += sizeof(T);
- return val;
-}
-
using namespace nvinfer1;
-#include "pluginsRT/ActivationLeakyRT.h"
-#include "pluginsRT/ActivationLogisticRT.h"
-#include "pluginsRT/ActivationReLUCeilingRT.h"
-#include "pluginsRT/ActivationMishRT.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/ResizeLayerRT.h"
-#include "pluginsRT/DeformableConvRT.h"
-#include "pluginsRT/FlattenConcatRT.h"
-#include "pluginsRT/ReshapeRT.h"
-#include "pluginsRT/MaxPoolingFixedSizeRT.h"
-
-class PluginFactory : IPluginFactory
-{
-public:
- YoloRT *yolos[16];
- int n_yolos;
-
- virtual IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength);
-};
@@ -69,12 +52,11 @@ public:
void* buffersRT[MAX_BUFFERS_RT];
dataDim_t buffersDIM[MAX_BUFFERS_RT];
int buf_input_idx, buf_output_idx;
-
+ bool builderActive = false;
dataDim_t input_dim, output_dim;
dnnType *output;
cudaStream_t stream;
- PluginFactory *pluginFactory;
NetworkRT(Network *net, const char *name);
virtual ~NetworkRT();
@@ -106,18 +88,26 @@ public:
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, Flatten *l);
- nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Reshape *l);
+ nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Flatten *l);
+ nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Reshape *l);
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Resize *l);
- nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Reorg *l);
- nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Region *l);
+ nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Reorg *l);
+ nvinfer1::IPluginV2Layer* 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::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Yolo *l);
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Upsample *l);
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, DeformConv2d *l);
+ nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input,Padding *l);
+ nvinfer1::ILayer* convert_layer(nvinfer1::ITensor* input,MulAdd *l);
+#if NV_TENSORRT_MAJOR > 5 && NV_TENSORRT_MAJOR < 8
bool serialize(const char *filename);
+#else
+ bool serialize(const char *filename,nvinfer1::IHostMemory *ptr);
+#endif
+
bool deserialize(const char *filename);
+ void destroy();
diff --git a/include/tkDNN/NetworkViz.h b/include/tkDNN/NetworkViz.h
index 2cf8009..ffdf361 100644
--- a/include/tkDNN/NetworkViz.h
+++ b/include/tkDNN/NetworkViz.h
@@ -6,7 +6,7 @@
namespace tk { namespace dnn {
cv::Mat vizFloat2colorMap(cv::Mat map, double min=0, double max=0, int classes=19);
-cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int img_h, int img_w, double min=0, double max=0, int classes=19);
+cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int img_h, int img_w, double min=0, double max=0, int classes=0);
cv::Mat vizLayer2Mat(tk::dnn::Network *net, int layer, int imgdim = 1000);
}}
diff --git a/include/tkDNN/SegmentationNN.h b/include/tkDNN/SegmentationNN.h
index b73ffa2..c39323c 100644
--- a/include/tkDNN/SegmentationNN.h
+++ b/include/tkDNN/SegmentationNN.h
@@ -4,7 +4,9 @@
#include
#include
#include
+#ifdef __linux__
#include
+#endif
#include
#include "utils.h"
@@ -181,6 +183,7 @@ class SegmentationNN {
checkCuda(cudaMemcpyAsync(mean_d, mean.data(), mean.size() * sizeof(float), cudaMemcpyHostToDevice, netRT->stream));
checkCuda(cudaMemcpyAsync(stddev_d, stddev.data(), stddev.size() * sizeof(float), cudaMemcpyHostToDevice, netRT->stream));
+ return true;
return true;
}
diff --git a/include/tkDNN/Yolo3Detection.h b/include/tkDNN/Yolo3Detection.h
index 100a720..5a29d9c 100644
--- a/include/tkDNN/Yolo3Detection.h
+++ b/include/tkDNN/Yolo3Detection.h
@@ -4,9 +4,9 @@
#include "opencv2/opencv.hpp"
#include "DetectionNN.h"
+#include "DarknetParser.h"
-namespace tk { namespace dnn {
-
+namespace tk { namespace dnn {
class Yolo3Detection : public DetectionNN
{
private:
@@ -19,12 +19,13 @@ private:
tk::dnn::Yolo* getYoloLayer(int n=0);
cv::Mat bgr_h;
+ std::vector noYolos;
public:
Yolo3Detection() {};
~Yolo3Detection() {};
- bool init(const std::string& tensor_path, const int n_classes=80, const int n_batches=1, const float conf_thresh=0.3);
+ bool init(const std::string& tensor_path,const std::string& cfg_path,const std::string& name_path,const int n_classes=80, const int n_batches=1, const float conf_thresh=0.3);
void preprocess(cv::Mat &frame, const int bi=0);
void postprocess(const int bi=0,const bool mAP=false);
};
diff --git a/include/tkDNN/demo_utils.h b/include/tkDNN/demo_utils.h
index 06a8970..c39e704 100644
--- a/include/tkDNN/demo_utils.h
+++ b/include/tkDNN/demo_utils.h
@@ -9,12 +9,13 @@
#ifdef __linux__
#include
+#endif
+
#include
#include
#include
#include
-#endif
void readCalibrationMatrix(const std::string& path, cv::Mat& calib_mat);
diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h
index d809129..4d5474b 100644
--- a/include/tkDNN/kernels.h
+++ b/include/tkDNN/kernels.h
@@ -48,4 +48,11 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
const int dst_dim, cudaStream_t stream = cudaStream_t(0));
void scalAdd(dnnType* dstData, int size, float alpha, float beta, int inc, cudaStream_t stream = cudaStream_t(0));
+
+void reflection_pad2d_out_forward(int32_t pad_h,int32_t pad_w,float *srcData,float *dstData,int32_t input_h,int32_t input_w,int32_t plane_dim,int32_t n_batch,cudaStream_t cudaStream = cudaStream_t(0));
+
+void constant_pad2d_forward(dnnType *srcData,dnnType *dstData,int32_t input_h,int32_t input_w,int32_t output_h,
+ int32_t output_w,int32_t c,int32_t n,int32_t padT,int32_t padL,dnnType constant,cudaStream_t cudaStream = cudaStream_t(0));
+
+
#endif //KERNELS_H
diff --git a/include/tkDNN/pluginsRT/ActivationLeakyRT.h b/include/tkDNN/pluginsRT/ActivationLeakyRT.h
index 330ed37..1d98a59 100644
--- a/include/tkDNN/pluginsRT/ActivationLeakyRT.h
+++ b/include/tkDNN/pluginsRT/ActivationLeakyRT.h
@@ -1,61 +1,88 @@
-#include
+#include "NvInfer.h"
#include "../kernels.h"
+#include
+#include
-class ActivationLeakyRT : public IPlugin {
+namespace nvinfer1 {
+ class ActivationLeakyRT : public IPluginV2 {
-public:
- ActivationLeakyRT(float s) {
- slope = s;
- }
+ public:
+ explicit ActivationLeakyRT(float s);
- ~ActivationLeakyRT(){
+ ActivationLeakyRT(const void *data, size_t length);
- }
+ ~ActivationLeakyRT();
- int getNbOutputs() const override {
- return 1;
- }
+ int getNbOutputs() const NOEXCEPT override;
- Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
- return inputs[0];
- }
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override;
- void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
- size = 1;
- for(int i=0; i 7
+ int enqueue(int batchSize, void const *const *inputs, void *const *outputs, void *workspace,
+ cudaStream_t stream) NOEXCEPT override;
+#elif NV_TENSORRT_MAJOR == 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
- virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+ size_t getSerializationSize() const NOEXCEPT override;
- activationLEAKYForward((dnnType*)reinterpret_cast(inputs[0]),
- reinterpret_cast(outputs[0]), batchSize*size, slope, stream);
- return 0;
- }
+ void serialize(void *buffer) const NOEXCEPT override;
+
+ bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override;
+
+ const char *getPluginType() const NOEXCEPT override;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ void destroy() NOEXCEPT override;
+
+ const char *getPluginNamespace() const NOEXCEPT override;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+ IPluginV2 *clone() const NOEXCEPT override;
+
+ int size;
+ float slope;
+
+ private:
+ std::string mPluginNamespace;
+ };
+
+ class ActivationLeakyRTPluginCreator : public IPluginCreator {
+ public:
+ ActivationLeakyRTPluginCreator();
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+ IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+ const char *getPluginName() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override;
+
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+ };
- virtual size_t getSerializationSize() override {
- return 1*sizeof(int) + 1*sizeof(float);
- }
-
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a=buf;
- tk::dnn::writeBUF(buf, size);
- assert(buf == a + getSerializationSize());
- }
-
- int size;
- float slope;
-};
+ REGISTER_TENSORRT_PLUGIN(ActivationLeakyRTPluginCreator);
+};
\ No newline at end of file
diff --git a/include/tkDNN/pluginsRT/ActivationLogisticRT.h b/include/tkDNN/pluginsRT/ActivationLogisticRT.h
index 063931f..d972752 100644
--- a/include/tkDNN/pluginsRT/ActivationLogisticRT.h
+++ b/include/tkDNN/pluginsRT/ActivationLogisticRT.h
@@ -1,60 +1,88 @@
#include
#include "../kernels.h"
+#include
+#include
+#include
-class ActivationLogisticRT : public IPlugin {
+namespace nvinfer1 {
-public:
- ActivationLogisticRT() {
+ class ActivationLogisticRT : public IPluginV2 {
+
+ public:
+ ActivationLogisticRT() ;
+
+ ActivationLogisticRT(const void *data, size_t length) ;
+
+ ~ActivationLogisticRT() ;
+
+ int getNbOutputs() const NOEXCEPT override ;
+
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
+
+ void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type,
+ PluginFormat format, int maxBatchSize) NOEXCEPT override ;
+
+ int initialize() NOEXCEPT override ;
+
+ void terminate() NOEXCEPT override ;
+
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override;
+
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+ cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR == 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
- }
+ size_t getSerializationSize() const NOEXCEPT override ;
- ~ActivationLogisticRT(){
+ void serialize(void *buffer) const NOEXCEPT override ;
- }
+ const char *getPluginType() const NOEXCEPT override ;
- int getNbOutputs() const override {
- return 1;
- }
+ const char *getPluginVersion() const NOEXCEPT override ;
- Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
- return inputs[0];
- }
+ void destroy() NOEXCEPT override ;
- void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
- size = 1;
- for(int i=0; i(inputs[0]),
- reinterpret_cast(outputs[0]), batchSize*size, stream);
- return 0;
- }
+ class ActivationLogisticRTPluginCreator : public IPluginCreator {
+ public:
+ ActivationLogisticRTPluginCreator() ;
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
- virtual size_t getSerializationSize() override {
- return 1*sizeof(int);
- }
+ IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer);
- tk::dnn::writeBUF(buf, size);
- }
+ const char *getPluginNamespace() const NOEXCEPT override ;
- int size;
-};
+ IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+ const char *getPluginName() const NOEXCEPT override ;
+
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+
+ REGISTER_TENSORRT_PLUGIN(ActivationLogisticRTPluginCreator);
+};
\ No newline at end of file
diff --git a/include/tkDNN/pluginsRT/ActivationMishRT.h b/include/tkDNN/pluginsRT/ActivationMishRT.h
index 5d660af..5b966cd 100644
--- a/include/tkDNN/pluginsRT/ActivationMishRT.h
+++ b/include/tkDNN/pluginsRT/ActivationMishRT.h
@@ -1,61 +1,82 @@
#include
#include "../kernels.h"
+#include
+#include
-class ActivationMishRT : public IPlugin {
+namespace nvinfer1 {
+ class ActivationMishRT : public IPluginV2 {
-public:
- ActivationMishRT() {
+ public:
+ ActivationMishRT() ;
+
+ ~ActivationMishRT() ;
+
+ ActivationMishRT(const void *data, size_t length) ;
- }
+ int getNbOutputs() const NOEXCEPT override ;
- ~ActivationMishRT(){
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
- }
+ void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type,
+ PluginFormat format, int maxBatchSize) NOEXCEPT override ;
- int getNbOutputs() const override {
- return 1;
- }
+ int initialize() NOEXCEPT override ;
- Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
- return inputs[0];
- }
+ void terminate() NOEXCEPT override ;
- void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
- size = 1;
- for(int i=0; i 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR == 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
- int initialize() override {
+ size_t getSerializationSize() const NOEXCEPT override ;
- return 0;
- }
+ void serialize(void *buffer) const NOEXCEPT override ;
- virtual void terminate() override {
- }
+ const char *getPluginType() const NOEXCEPT override ;
- virtual size_t getWorkspaceSize(int maxBatchSize) const override {
- return 0;
- }
+ const char *getPluginVersion() const NOEXCEPT override ;
- virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+ void destroy() NOEXCEPT override { delete this; }
- activationMishForward((dnnType*)reinterpret_cast(inputs[0]),
- reinterpret_cast(outputs[0]), batchSize*size, stream);
- return 0;
- }
+ bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+ const char *getPluginNamespace() const NOEXCEPT override ;
- virtual size_t getSerializationSize() override {
- return 1*sizeof(int);
- }
+ void setPluginNamespace(const char *plguinNamespace) NOEXCEPT override ;
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a=buf;
- tk::dnn::writeBUF(buf, size);
- assert(buf == a + getSerializationSize());
- }
+ IPluginV2 *clone() const NOEXCEPT override ;
- int size;
-};
+ int size;
+ private:
+ std::string mPluginNamespace;
+ };
+
+ class ActivationMishRTPluginCreator : public IPluginCreator {
+ public:
+ ActivationMishRTPluginCreator() ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+ IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+ const char *getPluginName() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+
+ REGISTER_TENSORRT_PLUGIN(ActivationMishRTPluginCreator);
+};
\ No newline at end of file
diff --git a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h
index 50ceb81..1830945 100644
--- a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h
+++ b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h
@@ -1,63 +1,81 @@
#include
#include "../kernels.h"
+#include
+#include
+#include
-class ActivationReLUCeiling : public IPlugin {
+namespace nvinfer1 {
+ class ActivationReLUCeiling : public IPluginV2 {
-public:
- ActivationReLUCeiling(const float ceiling) {
- this->ceiling = ceiling;
- }
+ public:
+ explicit ActivationReLUCeiling(const float ceiling) ;
- ~ActivationReLUCeiling(){
+ ~ActivationReLUCeiling() ;
- }
+ ActivationReLUCeiling(const void *data, size_t length) ;
- int getNbOutputs() const override {
- return 1;
- }
+ int getNbOutputs() const NOEXCEPT override ;
- Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
- return inputs[0];
- }
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
- void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
- size = 1;
- for(int i=0; i 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR == 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
- virtual size_t getWorkspaceSize(int maxBatchSize) const override {
- return 0;
- }
+ size_t getSerializationSize() const NOEXCEPT override ;
- virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+ void serialize(void *buffer) const NOEXCEPT override ;
- activationReLUCeilingForward((dnnType*)reinterpret_cast(inputs[0]),
- reinterpret_cast(outputs[0]), batchSize*size, ceiling, stream);
- return 0;
- }
+ IPluginV2 *clone() const NOEXCEPT override ;
+ bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
- virtual size_t getSerializationSize() override {
- return 1*sizeof(int) + 1*sizeof(float);
- }
+ void destroy() NOEXCEPT override ;
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a=buf;
- tk::dnn::writeBUF(buf, ceiling);
- tk::dnn::writeBUF(buf, size);
- assert(buf = a + getSerializationSize());
-
- }
+ const char *getPluginType() const NOEXCEPT override ;
- int size;
- float ceiling;
-};
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+ int size;
+ float ceiling;
+ private:
+ std::string mPluginNamespace;
+ };
+
+ class ActivationReLUCeilingPluginCreator : public IPluginCreator {
+ public:
+ ActivationReLUCeilingPluginCreator() ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+ IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+ const char *getPluginName() const NOEXCEPT override ;
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+ public:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+
+ REGISTER_TENSORRT_PLUGIN(ActivationReLUCeilingPluginCreator);
+};
\ No newline at end of file
diff --git a/include/tkDNN/pluginsRT/ConstantPaddingRT.h b/include/tkDNN/pluginsRT/ConstantPaddingRT.h
new file mode 100644
index 0000000..15f4c0d
--- /dev/null
+++ b/include/tkDNN/pluginsRT/ConstantPaddingRT.h
@@ -0,0 +1,109 @@
+//
+// Created by perseusdg on 1/7/22.
+//
+
+#ifndef _CONSTANTPADDINGRT_PLUGIN_H
+#define _CONSTANTPADDINGRT_PLUGIN_H
+
+#include
+#include
+#include
+#include
+#include
+
+namespace nvinfer1{
+ class ConstantPaddingRT : public IPluginV2Ext {
+ public:
+ ConstantPaddingRT(int32_t padH,int32_t padW,int32_t n,int32_t c,int32_t i_h,int32_t i_w,int32_t o_h,int32_t o_w,float constant);
+
+ ConstantPaddingRT(const void *data,size_t length);
+
+ ~ConstantPaddingRT();
+
+ int getNbOutputs() const NOEXCEPT override;
+
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
+
+ int initialize() NOEXCEPT override ;
+
+ void terminate() NOEXCEPT override ;
+
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+
+
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR <= 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+ size_t getSerializationSize() const NOEXCEPT override ;
+
+ void serialize(void *buffer) const NOEXCEPT override ;
+
+ void destroy() NOEXCEPT override ;
+
+ const char *getPluginType() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ IPluginV2Ext *clone() const NOEXCEPT override ;
+
+ DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+ void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+ bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+ bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+ void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+ int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+ bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+ int32_t maxBatchSize) NOEXCEPT override;
+
+ void detachFromContext() NOEXCEPT override;
+
+ bool supportsFormat (DataType type, PluginFormat format) const NOEXCEPT override;
+
+ int32_t i_h,i_w,o_h,o_w,n,c,padH,padW;
+ float constant;
+ private:
+ std::string mPluginNamespace;
+
+ };
+
+ class ConstantPaddingRTPluginCreator : public IPluginCreator {
+ public:
+ ConstantPaddingRTPluginCreator();
+
+ void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override;
+
+ const char *getPluginNamespace() const NOEXCEPT override;
+
+ IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+ IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+ const char *getPluginName() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+
+ };
+
+ REGISTER_TENSORRT_PLUGIN(ConstantPaddingRTPluginCreator);
+};
+
+
+#endif //TKDNN_CONSTANTPADDINGRT_H
diff --git a/include/tkDNN/pluginsRT/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h
index 711198f..9170e15 100644
--- a/include/tkDNN/pluginsRT/DeformableConvRT.h
+++ b/include/tkDNN/pluginsRT/DeformableConvRT.h
@@ -1,196 +1,137 @@
+#ifndef _DEFORMABLECONVRT_PLUGIN_H
+#define _DEFORMABLECONVRT_PLUGIN_H
+
+#include
+#include
#include
#include "../kernels.h"
+#include
+
+namespace nvinfer1 {
+ class DeformableConvRT : public IPluginV2Ext {
-class DeformableConvRT : public IPlugin {
+ public:
+ DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw,
+ int deformableGroup, int i_n, int i_c, int i_h, int i_w,
+ int o_n, int o_c, int o_h, int o_w,std::vector data_H,std::vector bias2_H,
+ std::vector ones_d1_h,std::vector ones_d2_h,std::vector offsetH,std::vector maskH,int height_ones,
+ int width_ones,int dim_ones);
+
+ ~DeformableConvRT();
+
+ DeformableConvRT(const void *data, size_t length) ;
+
+ int getNbOutputs() const NOEXCEPT override ;
+
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
+
+ int initialize() NOEXCEPT override ;
+
+ void terminate() NOEXCEPT override ;
+
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+ cudaStream_t stream) NOEXCEPT override;
+#elif NV_TENSORRT_MAJOR == 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+ size_t getSerializationSize() const NOEXCEPT override ;
+
+ void serialize(void *buffer) const NOEXCEPT override ;
+
+ void destroy() NOEXCEPT override ;
+
+ bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ const char *getPluginType() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ IPluginV2Ext *clone() const NOEXCEPT override ;
+
+ DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+ void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+ bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+ bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+ void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+ int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+ bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+ int32_t maxBatchSize) NOEXCEPT override;
+
+ void detachFromContext() NOEXCEPT override;
+ cublasStatus_t stat;
+ cublasHandle_t handle{nullptr};
+ int i_n, i_c, i_h, i_w;
+ int o_n, o_c, o_h, o_w;
+ int size;
+ int chunk_dim;
+ int kh, kw;
+ int sh, sw;
+ int ph, pw;
+ int deformableGroup;
+ int height_ones;
+ int width_ones;
+ int dim_ones;
-public:
- DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw,
- int deformableGroup, int i_n, int i_c, int i_h, int i_w,
- int o_n, int o_c, int o_h, int o_w,
- tk::dnn::DeformConv2d *deformable = nullptr) {
- this->chunk_dim = chunk_dim;
- this->kh = kh;
- this->kw = kw;
- this->sh = sh;
- this->sw = sw;
- this->ph = ph;
- this->pw = pw;
- this->deformableGroup = deformableGroup;
- this->i_n = i_n;
- this->i_c = i_c;
- this->i_h = i_h;
- this->i_w = i_w;
- this->o_n = o_n;
- this->o_c = o_c;
- this->o_h = o_h;
- this->o_w = o_w;
- height_ones = (i_h + 2 * ph - (1 * (kh - 1) + 1)) / sh + 1;
- width_ones = (i_w + 2 * pw - (1 * (kw - 1) + 1)) / sw + 1;
- dim_ones = i_c * kh * kw * 1 * height_ones * width_ones;
-
- checkCuda( cudaMalloc(&data_d, i_c * o_c * kh * kw * 1 * sizeof(dnnType)));
- checkCuda( cudaMalloc(&bias2_d, o_c*sizeof(dnnType)));
- checkCuda( cudaMalloc(&ones_d1, height_ones * width_ones * sizeof(dnnType)));
- checkCuda( cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType)));
- checkCuda( cudaMalloc(&mask, chunk_dim*sizeof(dnnType)));
- checkCuda( cudaMalloc(&ones_d2, dim_ones*sizeof(dnnType)));
- if(deformable != nullptr) {
- this->defRT = deformable;
- checkCuda( cudaMemcpy(data_d, deformable->data_d, sizeof(dnnType)*i_c * o_c * kh * kw * 1, cudaMemcpyDeviceToDevice) );
- checkCuda( cudaMemcpy(bias2_d, deformable->bias2_d, sizeof(dnnType)*o_c, cudaMemcpyDeviceToDevice) );
- checkCuda( cudaMemcpy(ones_d1, deformable->ones_d1, sizeof(dnnType)*height_ones*width_ones, cudaMemcpyDeviceToDevice) );
- checkCuda( cudaMemcpy(offset, deformable->offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToDevice) );
- checkCuda( cudaMemcpy(mask, deformable->mask, sizeof(dnnType)*chunk_dim, cudaMemcpyDeviceToDevice) );
- checkCuda( cudaMemcpy(ones_d2, deformable->ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToDevice) );
- }
- stat = cublasCreate(&handle);
- if (stat != CUBLAS_STATUS_SUCCESS)
- FatalError("CUBLAS initialization failed\n");
- }
-
- ~DeformableConvRT() {
- checkCuda( cudaFree(data_d) );
- checkCuda( cudaFree(bias2_d) );
- checkCuda( cudaFree(ones_d1) );
- checkCuda( cudaFree(offset) );
- checkCuda( cudaFree(mask) );
- checkCuda( cudaFree(ones_d2) );
- cublasDestroy(handle);
- }
-
- int getNbOutputs() const override {
- return 1;
- }
-
- Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
- return DimsCHW{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w};
- }
-
- void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { }
-
- 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(inputs[0]);
- dnnType *output_conv = (dnnType*)reinterpret_cast(inputs[1]);
-
- // split conv2d outputs into offset to mask
- for(int b=0; b(outputs[0]), ones_d2,
- kh, kw,
- sh, sw,
- ph, pw,
- 1, 1,
- deformableGroup, b,
- i_n, i_c, i_h, i_w,
- o_n, o_c, o_h, o_w,
- chunk_dim);
- }
- return 0;
- }
+ std::vector data_d_v;
+ std::vector bias2_d_v;
+ std::vector ones_d1_v;
+ std::vector offset_v;
+ std::vector mask_v;
+ std::vector ones_d2_v;
+ dnnType* data_d;
+ dnnType* bias2_d;
+ dnnType* ones_d1;
+ dnnType* offset;
+ dnnType* mask;
+ dnnType* ones_d2;
+ // dnnType *input_n;
+ // dnnType *offset_n;
+ // dnnType *mask_n;
+ // dnnType *output_n;
- virtual size_t getSerializationSize() override {
- return 16 * sizeof(int) + chunk_dim * 3 * sizeof(dnnType) + (i_c * o_c * kh * kw * 1 ) * sizeof(dnnType) +
- o_c * sizeof(dnnType) + height_ones * width_ones * sizeof(dnnType) + dim_ones * sizeof(dnnType);
- }
+ tk::dnn::DeformConv2d *defRT;
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a=buf;
- tk::dnn::writeBUF(buf, chunk_dim);
- tk::dnn::writeBUF(buf, kh);
- tk::dnn::writeBUF(buf, kw);
- tk::dnn::writeBUF(buf, sh);
- tk::dnn::writeBUF(buf, sw);
- tk::dnn::writeBUF(buf, ph);
- tk::dnn::writeBUF(buf, pw);
- tk::dnn::writeBUF(buf, deformableGroup);
- tk::dnn::writeBUF(buf, i_n);
- tk::dnn::writeBUF(buf, i_c);
- tk::dnn::writeBUF(buf, i_h);
- tk::dnn::writeBUF(buf, i_w);
- tk::dnn::writeBUF(buf, o_n);
- tk::dnn::writeBUF(buf, o_c);
- tk::dnn::writeBUF(buf, o_h);
- tk::dnn::writeBUF(buf, o_w);
- dnnType *aus = new dnnType[chunk_dim*2];
- checkCuda( cudaMemcpy(aus, offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToHost) );
- for(int i=0; i mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+ REGISTER_TENSORRT_PLUGIN(DeformableConvRTPluginCreator);
};
+#endif
diff --git a/include/tkDNN/pluginsRT/FlattenConcatRT.h b/include/tkDNN/pluginsRT/FlattenConcatRT.h
index 51aa1ab..f7ec495 100644
--- a/include/tkDNN/pluginsRT/FlattenConcatRT.h
+++ b/include/tkDNN/pluginsRT/FlattenConcatRT.h
@@ -1,81 +1,100 @@
+#ifndef _FLATTENCONCATRT_PLUGIN_H
+#define _FLATTENCONCATRT_PLUGIN_H
+
#include
+#include
+#include
+#include
+namespace nvinfer1 {
+ class FlattenConcatRT : public IPluginV2Ext {
-class FlattenConcatRT : public IPlugin {
+ public:
+ FlattenConcatRT(int c,int h,int w,int rows,int cols) ;
-public:
- FlattenConcatRT() {
- stat = cublasCreate(&handle);
- if (stat != CUBLAS_STATUS_SUCCESS) {
- printf ("CUBLAS initialization failed\n");
- return;
- }
- }
+ FlattenConcatRT(const void *data, size_t length) ;
- ~FlattenConcatRT(){
+ ~FlattenConcatRT() ;
- }
+ int getNbOutputs() const NOEXCEPT override ;
- int getNbOutputs() const override {
- return 1;
- }
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
- Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
- return DimsCHW{ inputs[0].d[0] * inputs[0].d[1] * inputs[0].d[2], 1, 1};
- }
+ int initialize() NOEXCEPT override ;
- void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
- assert(nbOutputs == 1 && nbInputs ==1);
- rows = inputDims[0].d[0];
- cols = inputDims[0].d[1] * inputDims[0].d[2];
- c = inputDims[0].d[0] * inputDims[0].d[1] * inputDims[0].d[2];
- h = 1;
- w = 1;
- }
+ void terminate() NOEXCEPT override ;
- int initialize() override {
- return 0;
- }
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
- virtual void terminate() override {
- checkERROR(cublasDestroy(handle));
- }
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR <= 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
- virtual size_t getWorkspaceSize(int maxBatchSize) const override {
- return 0;
- }
+ size_t getSerializationSize() const NOEXCEPT override ;
- virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
- dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]);
- dnnType *dstData = reinterpret_cast(outputs[0]);
- checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
+ void serialize(void *buffer) const NOEXCEPT override ;
- checkERROR( cublasSetStream(handle, stream) );
- for(int i=0; i(buffer),*a = buf;
- tk::dnn::writeBUF(buf, c);
- tk::dnn::writeBUF(buf, h);
- tk::dnn::writeBUF(buf, w);
- tk::dnn::writeBUF(buf, rows);
- tk::dnn::writeBUF(buf, cols);
- assert(buf == a + getSerializationSize());
- }
+ const char *getPluginNamespace() const NOEXCEPT override ;
- int c, h, w;
- int rows, cols;
- cublasStatus_t stat;
- cublasHandle_t handle;
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ IPluginV2Ext *clone() const NOEXCEPT override ;
+
+ DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+ void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+ bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+ bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+ void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+ int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+ bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+ int32_t maxBatchSize) NOEXCEPT override;
+
+ void detachFromContext() NOEXCEPT override;
+
+ bool supportsFormat (DataType type, PluginFormat format) const NOEXCEPT override;
+
+ int c, h, w;
+ int rows, cols;
+ cublasHandle_t handle{nullptr};
+ private:
+ std::string mPluginNamespace;
+ };
+
+ class FlattenConcatRTPluginCreator : public IPluginCreator {
+ public:
+ FlattenConcatRTPluginCreator() ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+ IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+ const char *getPluginName() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+
+ REGISTER_TENSORRT_PLUGIN(FlattenConcatRTPluginCreator);
};
+#endif
\ No newline at end of file
diff --git a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h
index 0899a34..95e15f4 100644
--- a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h
+++ b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h
@@ -1,75 +1,105 @@
#include
#include "../kernels.h"
-
-class MaxPoolFixedSizeRT : public IPlugin {
-
-public:
- MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding) {
- this->c = c;
- this->h = h;
- this->w = w;
- this->n = n;
- this->stride_H = strideH;
- this->stride_W = strideW;
- this->winSize = winSize;
- this->padding = padding;
- }
-
- ~MaxPoolFixedSizeRT(){
- }
-
- int getNbOutputs() const override {
- return 1;
- }
-
- Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
- return DimsCHW{this->c, this->h, this->w};
- }
-
- void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
- }
-
- 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 {
-
- //std::cout<n<<" "<c<<" "<h<<" "<w<<" "<stride_H<<" "<stride_W<<" "<winSize<<" "<padding<(inputs[0]);
- dnnType *dstData = reinterpret_cast(outputs[0]);
- MaxPoolingForward(srcData, dstData, batchSize, this->c, this->h, this->w, this->stride_H, this->stride_W, this->winSize, this->padding, stream);
- return 0;
- }
+#include
+#include
+#include
- virtual size_t getSerializationSize() override {
- return 8*sizeof(int);
- }
+namespace nvinfer1 {
+ class MaxPoolFixedSizeRT : public IPluginV2Ext {
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a=buf;
+ public:
+ MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding) ;
- tk::dnn::writeBUF(buf, this->c);
- tk::dnn::writeBUF(buf, this->h);
- tk::dnn::writeBUF(buf, this->w);
- tk::dnn::writeBUF(buf, this->n);
- tk::dnn::writeBUF(buf, this->stride_H);
- tk::dnn::writeBUF(buf, this->stride_W);
- tk::dnn::writeBUF(buf, this->winSize);
- tk::dnn::writeBUF(buf, this->padding);
- assert(buf == a + getSerializationSize());
- }
+ MaxPoolFixedSizeRT(const void *data, size_t length) ;
- int n, c, h, w;
- int stride_H, stride_W;
- int winSize;
- int padding;
+ ~MaxPoolFixedSizeRT() ;
+
+ int getNbOutputs() const NOEXCEPT override ;
+
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
+
+ int initialize() NOEXCEPT override ;
+
+ void terminate() NOEXCEPT override ;
+
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+ cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR <= 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+
+ size_t getSerializationSize() const NOEXCEPT override ;
+
+ void serialize(void *buffer) const NOEXCEPT override ;
+
+ void destroy() NOEXCEPT override ;
+
+ bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ const char *getPluginType() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ IPluginV2Ext *clone() const NOEXCEPT override ;
+
+ DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+ void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+ bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+ bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+ void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+ int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+ bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+ int32_t maxBatchSize) NOEXCEPT override;
+
+ void detachFromContext() NOEXCEPT override;
+
+
+ int n, c, h, w;
+ int stride_H, stride_W;
+ int winSize;
+ int padding;
+
+ private:
+ std::string mPluginNamespace;
+ };
+
+ class MaxPoolFixedSizeRTPluginCreator : public IPluginCreator {
+ public:
+ MaxPoolFixedSizeRTPluginCreator() ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+ IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+ const char *getPluginName() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+
+ };
+
+ REGISTER_TENSORRT_PLUGIN(MaxPoolFixedSizeRTPluginCreator);
};
diff --git a/include/tkDNN/pluginsRT/ReflectionPadding.h b/include/tkDNN/pluginsRT/ReflectionPadding.h
new file mode 100644
index 0000000..7b13710
--- /dev/null
+++ b/include/tkDNN/pluginsRT/ReflectionPadding.h
@@ -0,0 +1,101 @@
+#ifndef _REFLECTIONPADDINGRT_PLUGIN_H
+#define _REFLECTIONPADDINGRT_PLUGIN_H
+
+#include
+#include
+#include
+#include
+#include
+
+namespace nvinfer1{
+ class ReflectionPaddingRT : public IPluginV2Ext {
+ public:
+ ReflectionPaddingRT(int32_t padH,int32_t padW,int32_t input_h,int32_t input_w,int32_t output_h,int32_t output_w,int32_t c,int32_t n);
+
+ ReflectionPaddingRT(const void *data,size_t length);
+
+ ~ReflectionPaddingRT();
+
+ int getNbOutputs() const NOEXCEPT override;
+
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
+
+ int initialize() NOEXCEPT override ;
+
+ void terminate() NOEXCEPT override ;
+
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR <= 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+ size_t getSerializationSize() const NOEXCEPT override ;
+
+ void serialize(void *buffer) const NOEXCEPT override ;
+
+ void destroy() NOEXCEPT override ;
+
+ const char *getPluginType() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ IPluginV2Ext *clone() const NOEXCEPT override ;
+
+ DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+ void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+ bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+ bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+ void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+ int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+ bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+ int32_t maxBatchSize) NOEXCEPT override;
+
+ void detachFromContext() NOEXCEPT override;
+
+ bool supportsFormat (DataType type, PluginFormat format) const NOEXCEPT override;
+
+ int32_t padH,padW,input_h,input_w,output_h,output_w,n,c;
+ private:
+ std::string mPluginNamespace;
+
+ };
+
+ class ReflectionPaddingRTPluginCreator : public IPluginCreator {
+ public:
+ ReflectionPaddingRTPluginCreator();
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+ IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+ const char *getPluginName() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+
+ REGISTER_TENSORRT_PLUGIN(ReflectionPaddingRTPluginCreator);
+};
+#endif
+
diff --git a/include/tkDNN/pluginsRT/RegionRT.h b/include/tkDNN/pluginsRT/RegionRT.h
index 8487652..7f7157c 100644
--- a/include/tkDNN/pluginsRT/RegionRT.h
+++ b/include/tkDNN/pluginsRT/RegionRT.h
@@ -1,95 +1,110 @@
+#ifndef _REGIONRT_PLUGIN_H
+#define _REGIONRT_PLUGIN_H
#include
#include "../kernels.h"
+#include
+#include
+#include
-class RegionRT : public IPlugin {
+namespace nvinfer1 {
+ class RegionRT : public IPluginV2Ext {
-public:
- RegionRT(int classes, int coords, int num) {
+ public:
+ RegionRT(int classes, int coords, int num,int c,int h,int w);
- this->classes = classes;
- this->coords = coords;
- this->num = num;
- }
+ ~RegionRT() ;
- ~RegionRT(){
+ RegionRT(const void *data, size_t length) ;
- }
+ int getNbOutputs() const NOEXCEPT override ;
- int getNbOutputs() const override {
- return 1;
- }
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
- 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(inputs[0]);
- dnnType *dstData = reinterpret_cast(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);
- activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream);
-
- index = entry_index(b, n*w*h, coords);
- activationLOGISTICForward(srcData + index, dstData + index, w*h, stream);
- }
- }
-
- //softmax start
- int index = entry_index(0, 0, coords + 1);
- softmaxForward( srcData + index, classes, batchSize*num,
- (c*h*w)/num,
- w*h, 1, w*h, 1, dstData + index, stream);
-
- return 0;
- }
+ int initialize() NOEXCEPT override ;
- virtual size_t getSerializationSize() override {
- return 6*sizeof(int);
- }
+ void terminate() NOEXCEPT override ;
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a=buf;
- 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);
- assert(buf == a + getSerializationSize());
- }
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
- int c, h, w;
- int classes, coords, num;
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+ cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR == 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
- int entry_index(int batch, int location, int entry) {
- int n = location / (w*h);
- int loc = location % (w*h);
- return batch*c*h*w + n*w*h*(coords+classes+1) + entry*w*h + loc;
- }
+ size_t getSerializationSize() const NOEXCEPT override ;
+
+ void serialize(void *buffer) const NOEXCEPT override ;
+
+ const char *getPluginType() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ void destroy() NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+
+ DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+ void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+ bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+ bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+ void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+ int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+ bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+ int32_t maxBatchSize) NOEXCEPT override;
+
+ void detachFromContext() NOEXCEPT override;
+
+ IPluginV2Ext *clone() const NOEXCEPT override ;
+ int c, h, w;
+ int classes, coords, num;
+
+ int entry_index(int batch, int location, int entry) {
+ int n = location / (w * h);
+ int loc = location % (w * h);
+ return batch * c * h * w + n * w * h * (coords + classes + 1) + entry * w * h + loc;
+ }
+
+ private:
+ std::string mPluginNamespace;
+ };
+
+ class RegionRTPluginCreator : public IPluginCreator {
+ public:
+ RegionRTPluginCreator();
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+ IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+ const char *getPluginName() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+
+ REGISTER_TENSORRT_PLUGIN(RegionRTPluginCreator);
};
+
+#endif
diff --git a/include/tkDNN/pluginsRT/ReorgRT.h b/include/tkDNN/pluginsRT/ReorgRT.h
index c1b529a..be163a5 100644
--- a/include/tkDNN/pluginsRT/ReorgRT.h
+++ b/include/tkDNN/pluginsRT/ReorgRT.h
@@ -1,64 +1,98 @@
#include
#include "../kernels.h"
+#include
+#include
-class ReorgRT : public IPlugin {
+namespace nvinfer1 {
+ class ReorgRT : public IPluginV2Ext {
-public:
- ReorgRT(int stride) {
- this->stride = stride;
- }
+ public:
+ ReorgRT(int stride,int c,int h,int w);
- ~ReorgRT(){
+ ~ReorgRT();
- }
+ ReorgRT(const void *data, size_t length);
- int getNbOutputs() const override {
- return 1;
- }
+ int getNbOutputs() const NOEXCEPT override;
- 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};
- }
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override;
- 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() NOEXCEPT override;
- int initialize() override {
+ void terminate() NOEXCEPT override;
- return 0;
- }
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override;
- 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(inputs[0]),
- reinterpret_cast(outputs[0]),
- batchSize, c, h, w, stride, stream);
- return 0;
- }
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+ cudaStream_t stream) NOEXCEPT override;
+#elif NV_TENSORRT_MAJOR == 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
- virtual size_t getSerializationSize() override {
- return 4*sizeof(int);
- }
+ size_t getSerializationSize() const NOEXCEPT override;
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a=buf;
- tk::dnn::writeBUF(buf, stride);
- tk::dnn::writeBUF(buf, c);
- tk::dnn::writeBUF(buf, h);
- tk::dnn::writeBUF(buf, w);
- assert(buf == a + getSerializationSize());
- }
+ void serialize(void *buffer) const NOEXCEPT override;
- int c, h, w, stride;
+ bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override;
+
+ const char *getPluginType() const NOEXCEPT override;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ void destroy() NOEXCEPT override;
+
+ const char *getPluginNamespace() const NOEXCEPT override;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+ IPluginV2Ext *clone() const NOEXCEPT override;
+
+ DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+ void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+ bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+ bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+ void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+ int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+ bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+ int32_t maxBatchSize) NOEXCEPT override;
+
+ void detachFromContext() NOEXCEPT override;
+
+ int c, h, w, stride;
+ private:
+ std::string mPluginNamespace;
+ };
+
+ class ReorgRTPluginCreator : public IPluginCreator {
+ public:
+ ReorgRTPluginCreator();
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+ const char *getPluginNamespace() const NOEXCEPT override;
+
+ IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override;
+
+ IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override;
+
+ const char *getPluginName() const NOEXCEPT override;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override;
+
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+
+ REGISTER_TENSORRT_PLUGIN(ReorgRTPluginCreator);
};
+
diff --git a/include/tkDNN/pluginsRT/ReshapeRT.h b/include/tkDNN/pluginsRT/ReshapeRT.h
index 37017c7..e56c79c 100644
--- a/include/tkDNN/pluginsRT/ReshapeRT.h
+++ b/include/tkDNN/pluginsRT/ReshapeRT.h
@@ -1,62 +1,102 @@
+#ifndef _RESHAPERT_PLUGIN_H
+#define _RESHAPERT_PLUGIN_H
+
#include
-
-class ReshapeRT : public IPlugin {
-
-public:
- ReshapeRT(dataDim_t new_dim) {
- n = new_dim.n;
- c = new_dim.c;
- h = new_dim.h;
- w = new_dim.w;
- }
-
- ~ReshapeRT(){
-
- }
-
- int getNbOutputs() const override {
- return 1;
- }
-
- Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
- return DimsCHW{ c,h,w};
- }
-
- void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
- }
-
- 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(inputs[0]);
- dnnType *dstData = reinterpret_cast(outputs[0]);
-
- checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
- return 0;
- }
+#include
+#include
+#include
+using namespace tk::dnn;
- virtual size_t getSerializationSize() override {
- return 4*sizeof(int);
- }
+namespace nvinfer1 {
+ class ReshapeRT : public IPluginV2Ext {
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a = buf;
- tk::dnn::writeBUF(buf, n);
- tk::dnn::writeBUF(buf, c);
- tk::dnn::writeBUF(buf, h);
- tk::dnn::writeBUF(buf, w);
- assert(buf == a + getSerializationSize());
- }
+ public:
+ ReshapeRT(int n,int c,int h,int w) ;
- int n, c, h, w;
+ ReshapeRT(const void *data, size_t length) ;
+
+ ~ReshapeRT() ;
+
+ int getNbOutputs() const NOEXCEPT override ;
+
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
+
+
+ int initialize() NOEXCEPT override ;
+
+ void terminate() NOEXCEPT override ;
+
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR == 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+ size_t getSerializationSize() const NOEXCEPT override ;
+
+ void serialize(void *buffer) const NOEXCEPT override ;
+
+ bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+
+ const char *getPluginType() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ void destroy() NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ IPluginV2Ext *clone() const NOEXCEPT override ;
+
+ DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+ void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+ bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+ bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+ void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+ int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+ bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+ int32_t maxBatchSize) NOEXCEPT override;
+
+ void detachFromContext() NOEXCEPT override;
+
+ int n, c, h, w;
+ private:
+ std::string mPluginNamespace;
+ };
+
+ class ReshapeRTPluginCreator : public IPluginCreator {
+ public:
+ ReshapeRTPluginCreator() ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+ IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+ const char *getPluginName() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+
+ REGISTER_TENSORRT_PLUGIN(ReshapeRTPluginCreator);
};
+#endif
\ No newline at end of file
diff --git a/include/tkDNN/pluginsRT/ResizeLayerRT.h b/include/tkDNN/pluginsRT/ResizeLayerRT.h
index cde52bf..750057e 100644
--- a/include/tkDNN/pluginsRT/ResizeLayerRT.h
+++ b/include/tkDNN/pluginsRT/ResizeLayerRT.h
@@ -1,68 +1,104 @@
#include
#include "../kernels.h"
+#include
+#include
+#include
-class ResizeLayerRT : public IPlugin {
+namespace nvinfer1 {
-public:
- ResizeLayerRT(int c, int h, int w) {
- o_c = c;
- o_h = h;
- o_w = w;
- }
+ class ResizeLayerRT : public IPluginV2Ext {
- ~ResizeLayerRT(){
- }
+ public:
+ ResizeLayerRT(int oc, int oh, int ow,int ic,int ih,int iw) ;
- int getNbOutputs() const override {
- return 1;
- }
+ ResizeLayerRT(const void *data, size_t length) ;
- Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
- return DimsCHW{o_c, o_h, o_w};
- }
+ ~ResizeLayerRT() ;
- void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
- i_c = inputDims[0].d[0];
- i_h = inputDims[0].d[1];
- i_w = inputDims[0].d[2];
- }
+ int getNbOutputs() const NOEXCEPT override ;
- int initialize() override {
- return 0;
- }
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
- virtual void terminate() override {
- }
+ int initialize() NOEXCEPT override ;
- virtual size_t getWorkspaceSize(int maxBatchSize) const override {
- return 0;
- }
+ void terminate() NOEXCEPT override ;
- virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
- // printf("%d %d %d %d %d %d\n", i_c, i_w, i_h, o_c, o_w, o_h);
- resizeForward((dnnType*)reinterpret_cast(inputs[0]),
- reinterpret_cast(outputs[0]),
- batchSize, i_c, i_h, i_w, o_c, o_h, o_w, stream);
- return 0;
- }
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+ cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR <= 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+ size_t getSerializationSize() const NOEXCEPT override ;
+
+ void serialize(void *buffer) const NOEXCEPT override ;
+
+ bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+
+ const char *getPluginType() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ void destroy() NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ IPluginV2Ext *clone() const NOEXCEPT override ;
+
+ DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+ void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+ bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+ bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+ void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+ int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+ bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+ int32_t maxBatchSize) NOEXCEPT override;
+
+ void detachFromContext() NOEXCEPT override;
+
+ int i_c, i_h, i_w, o_c, o_h, o_w;
+
+ private:
+ std::string mPluginNamespace;
+ };
+
+ class ResizeLayerRTPluginCreator : public IPluginCreator {
+ public:
+ ResizeLayerRTPluginCreator() ;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+ const char *getPluginNamespace() const NOEXCEPT override ;
+
+ IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+ IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+ const char *getPluginName() const NOEXCEPT override ;
+
+ const char *getPluginVersion() const NOEXCEPT override ;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
- virtual size_t getSerializationSize() override {
- return 6*sizeof(int);
- }
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a=buf;
- tk::dnn::writeBUF(buf, o_c);
- tk::dnn::writeBUF(buf, o_h);
- tk::dnn::writeBUF(buf, o_w);
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
- tk::dnn::writeBUF(buf, i_c);
- tk::dnn::writeBUF(buf, i_h);
- tk::dnn::writeBUF(buf, i_w);
- assert(buf == a + getSerializationSize());
- }
+ };
- int i_c, i_h, i_w, o_c, o_h, o_w;
+ REGISTER_TENSORRT_PLUGIN(ResizeLayerRTPluginCreator);
};
+
diff --git a/include/tkDNN/pluginsRT/RouteRT.h b/include/tkDNN/pluginsRT/RouteRT.h
index 5a8c170..499b9da 100644
--- a/include/tkDNN/pluginsRT/RouteRT.h
+++ b/include/tkDNN/pluginsRT/RouteRT.h
@@ -1,96 +1,90 @@
#include
#include "../kernels.h"
+#include
+#include
-class RouteRT : public IPlugin {
+namespace nvinfer1 {
+ class RouteRT : public IPluginV2 {
- /**
- THIS IS NOT USED ANYMORE
- */
+ /**
+ THIS IS NOT USED ANYMORE
+ */
-public:
- RouteRT(int groups, int group_id) {
- this->groups = groups;
- this->group_id = group_id;
- }
+ public:
+ RouteRT(int groups, int group_id) ;
- ~RouteRT(){
+ ~RouteRT() ;
- }
+ RouteRT(const void *data, size_t length) ;
- int getNbOutputs() const override {
- return 1;
- }
+ int getNbOutputs() const NOEXCEPT override ;
- Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
- int out_c = 0;
- for(int i=0; i 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR == 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
- virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
-
- dnnType *dstData = reinterpret_cast(outputs[0]);
+ size_t getSerializationSize() const NOEXCEPT override ;
- for(int b=0; b(inputs[i]);
- int in_dim = c_in[i]*h*w;
- int part_in_dim = in_dim / this->groups;
- checkCuda( cudaMemcpyAsync(dstData + b*c*w*h + offset, input + b*c*w*h*groups + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) );
- offset += part_in_dim;
- }
- }
+ void serialize(void *buffer) const NOEXCEPT override ;
- return 0;
- }
+ const char *getPluginType() const NOEXCEPT override ;
+ const char *getPluginVersion() const NOEXCEPT override ;
- virtual size_t getSerializationSize() override {
- return (6+MAX_INPUTS)*sizeof(int);
- }
+ void destroy() NOEXCEPT override ;
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a=buf;
- tk::dnn::writeBUF(buf, groups);
- tk::dnn::writeBUF(buf, group_id);
- tk::dnn::writeBUF(buf, in);
- for(int i=0; i mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+
+ REGISTER_TENSORRT_PLUGIN(RouteRTPluginCreator);
};
diff --git a/include/tkDNN/pluginsRT/ShortcutRT.h b/include/tkDNN/pluginsRT/ShortcutRT.h
index 04091ac..0c01b9d 100644
--- a/include/tkDNN/pluginsRT/ShortcutRT.h
+++ b/include/tkDNN/pluginsRT/ShortcutRT.h
@@ -1,77 +1,109 @@
+#ifndef _SHORTCUTRT_PLUGIN_H
+#define _SHORTCUTRT_PLUGIN_H
+
#include
#include "../kernels.h"
-
-class ShortcutRT : public IPlugin {
-
-public:
- ShortcutRT(tk::dnn::dataDim_t bdim, bool mul) {
- this->bc = bdim.c;
- this->bh = bdim.h;
- this->bw = bdim.w;
- this->mul = mul;
- }
-
- ~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(inputs[0]);
- dnnType *srcDataBack = (dnnType*)reinterpret_cast(inputs[1]);
- dnnType *dstData = reinterpret_cast(outputs[0]);
-
- checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
- shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, bc, bh, bw, 1, mul, stream);
-
- return 0;
- }
+#include
+#include
+#include
- virtual size_t getSerializationSize() override {
- return 6*sizeof(int) + sizeof(bool);
- }
+namespace nvinfer1 {
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a=buf;
- tk::dnn::writeBUF(buf, bc);
- tk::dnn::writeBUF(buf, bh);
- tk::dnn::writeBUF(buf, bw);
- tk::dnn::writeBUF(buf, mul);
- tk::dnn::writeBUF(buf, c);
- tk::dnn::writeBUF(buf, h);
- tk::dnn::writeBUF(buf, w);
- assert(buf == a + getSerializationSize());
-
- }
+ class ShortcutRT : public IPluginV2Ext {
+
+ public:
+ ShortcutRT(int bc,int bh,int bw,int c,int h,int w ,bool mul);
+
+ ~ShortcutRT();
+
+ ShortcutRT(const void *data, size_t length);
+
+ int getNbOutputs() const NOEXCEPT override;
+
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override;
+
+ void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims, int32_t nbOutputs,
+ DataType const *inputTypes, DataType const *outputTypes, bool const *inputIsBroadcast,
+ bool const *outputIsBroadcast, PluginFormat floatFormat, int32_t maxBatchSize) NOEXCEPT override;
+
+ bool isOutputBroadcastAcrossBatch (int32_t outputIndex, bool const *inputIsBroadcasted, int32_t nbInputs) const NOEXCEPT override;
+
+ bool canBroadcastInputAcrossBatch (int32_t inputIndex) const NOEXCEPT override;
+
+ void attachToContext (cudnnContext *, cublasContext *, IGpuAllocator *) NOEXCEPT override;
+
+ void detachFromContext () NOEXCEPT override;
+
+ DataType getOutputDataType(int32_t index, nvinfer1::DataType const *inputTypes, int32_t nbInputs) const NOEXCEPT override;
+
+ int initialize() NOEXCEPT override;
+
+ void terminate() NOEXCEPT override;
+
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override;
+
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+ cudaStream_t stream) NOEXCEPT override;
+#elif NV_TENSORRT_MAJOR == 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+
+ size_t getSerializationSize() const NOEXCEPT override;
+
+ void serialize(void *buffer) const NOEXCEPT override;
+
+ bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override;
+
+ const char *getPluginType() const NOEXCEPT override;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ void destroy() NOEXCEPT override;
+
+ const char *getPluginNamespace() const NOEXCEPT override;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+ IPluginV2Ext *clone() const NOEXCEPT override;
+
+ int c, h, w;
+ int bc, bh, bw,bl;
+ bool mul;
+ tk::dnn::dataDim_t bDim;
+ private:
+ std::string mPluginNamespace;
+ };
+
+
+ class ShortcutRTPluginCreator : public IPluginCreator {
+ public:
+ ShortcutRTPluginCreator();
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+ const char *getPluginNamespace() const NOEXCEPT override;
+
+ IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override;
+
+ IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override;
+
+ const char *getPluginName() const NOEXCEPT override;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override;
+
+ public:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+
+ REGISTER_TENSORRT_PLUGIN(ShortcutRTPluginCreator);
- int c, h, w;
- int bc, bh, bw;
- bool mul;
};
+
+#endif
\ No newline at end of file
diff --git a/include/tkDNN/pluginsRT/UpsampleRT.h b/include/tkDNN/pluginsRT/UpsampleRT.h
index a11d7b4..4379ee7 100644
--- a/include/tkDNN/pluginsRT/UpsampleRT.h
+++ b/include/tkDNN/pluginsRT/UpsampleRT.h
@@ -1,66 +1,103 @@
+#ifndef _UPSAMPLERT_PLUGIN_H
+#define _UPSAMPLERT_PLUGIN_H
+
#include
#include "../kernels.h"
+#include
+#include
-class UpsampleRT : public IPlugin {
+namespace nvinfer1 {
-public:
- UpsampleRT(int stride) {
- this->stride = stride;
- }
+ class UpsampleRT : public IPluginV2Ext {
- ~UpsampleRT(){
+ public:
+ UpsampleRT(int stride,int c,int h,int w);
- }
+ UpsampleRT(const void *data, size_t length);
- int getNbOutputs() const override {
- return 1;
- }
+ ~UpsampleRT();
- 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);
- }
+ int getNbOutputs() const NOEXCEPT override;
- 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];
- }
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override;
- int initialize() override {
+ int initialize() NOEXCEPT override;
- return 0;
- }
+ void terminate() NOEXCEPT override;
- virtual void terminate() override {
- }
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT 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(inputs[0]);
- dnnType *dstData = reinterpret_cast(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;
- }
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+ cudaStream_t stream) NOEXCEPT override;
+#elif NV_TENSORRT_MAJOR == 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
- virtual size_t getSerializationSize() override {
- return 4*sizeof(int);
- }
+ size_t getSerializationSize() const NOEXCEPT override;
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a=buf;
- tk::dnn::writeBUF(buf, stride);
- tk::dnn::writeBUF(buf, c);
- tk::dnn::writeBUF(buf, h);
- tk::dnn::writeBUF(buf, w);
- assert(buf == a + getSerializationSize());
- }
+ void serialize(void *buffer) const NOEXCEPT override;
- int c, h, w, stride;
-};
+ bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override;
+
+ const char *getPluginType() const NOEXCEPT override;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ void destroy() NOEXCEPT override;
+
+ const char *getPluginNamespace() const NOEXCEPT override;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+ IPluginV2Ext *clone() const NOEXCEPT override ;
+
+ bool isOutputBroadcastAcrossBatch (int32_t outputIndex, bool const *inputIsBroadcasted, int32_t nbInputs) const NOEXCEPT override;
+
+ bool canBroadcastInputAcrossBatch (int32_t inputIndex) const NOEXCEPT override;
+
+ void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims, int32_t nbOutputs,
+ DataType const *inputTypes, DataType const *outputTypes, bool const *inputIsBroadcast,
+ bool const *outputIsBroadcast, PluginFormat floatFormat, int32_t maxBatchSize) NOEXCEPT override;
+
+ void attachToContext (cudnnContext *, cublasContext *, IGpuAllocator *) NOEXCEPT override;
+
+ void detachFromContext () NOEXCEPT override;
+
+ DataType getOutputDataType (int32_t index, nvinfer1::DataType const *inputTypes, int32_t nbInputs) const NOEXCEPT override;
+
+
+ int c, h, w, stride;
+ private:
+ std::string mPluginNamespace;
+ };
+
+ class UpsampleRTPluginCreator : public IPluginCreator {
+ public:
+ UpsampleRTPluginCreator();
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+ const char *getPluginNamespace() const NOEXCEPT override;
+
+ IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override;
+
+ IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override;
+
+ const char *getPluginName() const NOEXCEPT override;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override;
+
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+
+ REGISTER_TENSORRT_PLUGIN(UpsampleRTPluginCreator);
+ };
+
+#endif
\ No newline at end of file
diff --git a/include/tkDNN/pluginsRT/YoloRT.h b/include/tkDNN/pluginsRT/YoloRT.h
index 5ffe39c..352a0cb 100644
--- a/include/tkDNN/pluginsRT/YoloRT.h
+++ b/include/tkDNN/pluginsRT/YoloRT.h
@@ -1,143 +1,124 @@
+#ifndef _YOLORT_PLUGIN_H
+#define _YOLORT_PLUGIN_H
+
#include
+#include
#include "../kernels.h"
+#include
+#include
#define YOLORT_CLASSNAME_W 256
-class YoloRT : public IPlugin {
+namespace nvinfer1 {
+ class YoloRT : public IPluginV2Ext {
+
+ public:
+ YoloRT(int classes, int num,int c,int h,int w, int n_masks = 3, float scale_xy = 1,
+ float nms_thresh = 0.45, int nms_kind = 0, int new_coords = 0);
+
+ YoloRT(const void *data, size_t length);
+
+ ~YoloRT();
+ int getNbOutputs() const NOEXCEPT override;
-public:
- YoloRT(int classes, int num, tk::dnn::Yolo *yolo = nullptr, int n_masks=3, float scale_xy=1, float nms_thresh=0.45, int nms_kind=0, int new_coords=0) {
+ Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override;
- this->classes = classes;
- this->num = num;
- this->n_masks = n_masks;
- this->scaleXY = scale_xy;
- this->nms_thresh = nms_thresh;
- this->nms_kind = nms_kind;
- this->new_coords = new_coords;
+ int initialize() NOEXCEPT override;
- mask = new dnnType[n_masks];
- bias = new dnnType[num*n_masks*2];
- if(yolo != nullptr) {
- memcpy(mask, yolo->mask_h, sizeof(dnnType)*n_masks);
- memcpy(bias, yolo->bias_h, sizeof(dnnType)*num*n_masks*2);
- classesNames = yolo->classesNames;
- }
- }
+ void terminate() NOEXCEPT override;
- ~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(inputs[0]);
- dnnType *dstData = reinterpret_cast(outputs[0]);
-
- checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
+ size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override;
- for (int b = 0; b < batchSize; ++b){
- for(int n = 0; n < n_masks; ++n){
- int index = entry_index(b, n*w*h, 0);
- if (new_coords == 1){
- if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
- }
- else{
- activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); //x,y
+#if NV_TENSORRT_MAJOR > 7
+ int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+ cudaStream_t stream) NOEXCEPT override;
+#elif NV_TENSORRT_MAJOR == 7
+ int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
- if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
- index = entry_index(b, n*w*h, 4);
- activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream);
- }
- }
+ size_t getSerializationSize() const NOEXCEPT override;
+
+ bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override;
+
+ void serialize(void *buffer) const NOEXCEPT override;
+
+ const char *getPluginType() const NOEXCEPT override;
+
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ void destroy() NOEXCEPT override;
+
+ const char *getPluginNamespace() const NOEXCEPT override;
+
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+ IPluginV2Ext *clone() const NOEXCEPT override;
+
+ DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+ void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+ bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+ bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+ void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+ int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+ bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+ int32_t maxBatchSize) NOEXCEPT override;
+
+ void detachFromContext() NOEXCEPT override;
+
+
+ int c, h, w;
+ int classes, num, n_masks;
+ float scaleXY;
+ float nms_thresh;
+ int nms_kind;
+ int new_coords;
+ int NUM = 0;
+ std::vector classesNames;
+
+
+ int entry_index(int batch, int location, int entry) {
+ int n = location / (w * h);
+ int loc = location % (w * h);
+ return batch * c * h * w + n * w * h * (4 + classes + 1) + entry * w * h + loc;
}
- //std::cout<<"YOLO END\n";
- return 0;
- }
+ private:
+ std::string mPluginNamespace;
+ };
- virtual size_t getSerializationSize() override {
- return 8*sizeof(int) + 2*sizeof(float)+ n_masks*sizeof(dnnType) + num*n_masks*2*sizeof(dnnType) + YOLORT_CLASSNAME_W*classes*sizeof(char);
- }
+ class YoloRTPluginCreator : public IPluginCreator {
+ public:
+ YoloRTPluginCreator();
- virtual void serialize(void* buffer) override {
- char *buf = reinterpret_cast(buffer),*a=buf;
- tk::dnn::writeBUF(buf, classes); //std::cout << "Classes :" << classes << std::endl;
- tk::dnn::writeBUF(buf, num); //std::cout << "Num : " << num << std::endl;
- tk::dnn::writeBUF(buf, n_masks); //std::cout << "N_Masks" << n_masks << std::endl;
- tk::dnn::writeBUF(buf, scaleXY); //std::cout << "ScaleXY :" << scaleXY << std::endl;
- tk::dnn::writeBUF(buf, nms_thresh); //std::cout << "nms_thresh :" << nms_thresh << std::endl;
- tk::dnn::writeBUF(buf, nms_kind); //std::cout << "nms_kind : " << nms_kind << std::endl;
- tk::dnn::writeBUF(buf, new_coords); //std::cout << "new_coords : " << new_coords << std::endl;
- tk::dnn::writeBUF(buf, c); //std::cout << "C : " << c << std::endl;
- tk::dnn::writeBUF(buf, h); //std::cout << "H : " << h << std::endl;
- tk::dnn::writeBUF(buf, w); //std::cout << "C : " << c << std::endl;
- for (int i = 0; i < n_masks; i++)
- {
- tk::dnn::writeBUF(buf, mask[i]); //std::cout << "mask[i] : " << mask[i] << std::endl;
- }
- for (int i = 0; i < n_masks * 2 * num; i++)
- {
- tk::dnn::writeBUF(buf, bias[i]); //std::cout << "bias[i] : " << bias[i] << std::endl;
- }
+ void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
- // save classes names
- for(int i=0; i classesNames;
+ IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override;
- dnnType *mask;
- dnnType *bias;
+ IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override;
- int entry_index(int batch, int location, int entry) {
- int n = location / (w*h);
- int loc = location % (w*h);
- return batch*c*h*w + n*w*h*(4+classes+1) + entry*w*h + loc;
- }
+ const char *getPluginName() const NOEXCEPT override;
+ const char *getPluginVersion() const NOEXCEPT override;
+
+ const PluginFieldCollection *getFieldNames() NOEXCEPT override;
+
+ private:
+ static PluginFieldCollection mFC;
+ static std::vector mPluginAttributes;
+ std::string mPluginNamespace;
+ };
+
+ REGISTER_TENSORRT_PLUGIN(YoloRTPluginCreator);
};
+#endif
\ No newline at end of file
diff --git a/include/tkDNN/utils.h b/include/tkDNN/utils.h
index 4965a7a..055ea67 100644
--- a/include/tkDNN/utils.h
+++ b/include/tkDNN/utils.h
@@ -6,16 +6,18 @@
#include