diff --git a/.gitignore b/.gitignore index 683834c..22f7f94 100644 --- a/.gitignore +++ b/.gitignore @@ -21,3 +21,4 @@ scripts/COCO_val2017/* scripts/COCO_val2017.zip scripts/all_labels.txt /cmake/cuda_script +/cmake-build-debug/ diff --git a/CMakeLists.txt b/CMakeLists.txt index 7cc9e33..d919b46 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -85,12 +85,14 @@ endif() 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" "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} ${CUDA_LIBRARIES} ${CUDNN_LIBRARIES}) +target_link_libraries(kernels ${CUDA_CUBLAS_LIBRARIES} ${CUDA_LIBRARIES} ${CUDNN_LIBRARIES} yaml-cpp) @@ -120,7 +122,6 @@ endif() # endif() # gives problems in cross-compiling, probably malformed cmake config -find_package(yaml-cpp REQUIRED) #------------------------------------------------------------------------------- # Build Libraries diff --git a/README.md b/README.md index d07ac42..af7c9c5 100644 --- a/README.md +++ b/README.md @@ -17,9 +17,15 @@ If you use tkDNN in your research, please cite the [following paper](https://iee } ``` -### What's new (November 2021) -- [x] Support to sematic segmentation on cuda 11+ [README](docs/README_seg.md) -- [x] Support to TensorRT8 +### What's new +#### 20 July 2021 +- [x] Support to sematic segmentation [README](docs/README_seg.md) +- [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. + +TensorRT8 (and therefore Jetpack 4.6) is currently supported only on the branch tensort8 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 @@ -81,7 +87,7 @@ Results for COCO val 2017 (5k images), on RTX 2080Ti, with conf threshold=0.001 ## Dependencies This branch works on every NVIDIA GPU that supports the following (latest tested) dependencies: -* CUDA 11.3 (or >= 10.2) [the segmentation only works with CUDA 10 for now] +* 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) diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index 967c996..6de7214 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -18,64 +18,59 @@ void sig_handler(int signo) { int main(int argc, char *argv[]) { - std::cout<<"detection\n"; signal(SIGINT, sig_handler); +#ifdef __linux__ + std::string config_file = "../demo/demoConfig.yaml"; +#elif _WIN32 + std::string config_file = "..\\..\\..\\demo\\demoConfig.yaml"; +#endif - std::string net = "yolo4tiny_fp32.rt"; - #ifdef __linux__ - std::string cfgPath = "../tests/darknet/cfg/yolo4tiny.cfg"; - #elif _WIN32 - std::string cfgPath = "..\\tests\\darknet\\cfg\\yolo4tiny.cfg"; - #endif - - #ifdef __linux__ - std::string namePath = "../tests/darknet/names/coco.names"; - #elif _WIN32 - std::string namePath = "..\\tests\\darknet\\names\\coco.names"; - #endif - - 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 - - char ntype = 'y'; - if(argc > 2) - input = argv[2]; - if(argc > 3) - ntype = argv[3][0]; - int n_classes = 80; - if(argc > 4) - n_classes = atoi(argv[4]); - if(argc > 5) - cfgPath = argv[5]; - if(argc > 6) - namePath = argv[6]; - int n_batch = 1; - if(argc > 7) - n_batch = atoi(argv[7]); - bool show = true; - if(argc > 8) - show = atoi(argv[8]); - float conf_thresh=0.3; - if(argc >= 9) - conf_thresh = atof(argv[9]); - - if(n_batch < 1 || n_batch > 64) - FatalError("Batch dim not supported"); - - if(!show) - SAVE_RESULT = true; - - if(ntype == 'c' || ntype == 'm'){ - cfgPath = ""; - namePath = ""; - + if(argc > 1){ + config_file = argv[1]; } + + YAML::Node conf = YAMLloadConf(config_file); + if(!conf){ + FatalError("Problem with config file"); + } + + + std::string net = YAMLgetConf(conf,"net","yolo4tiny_fp32.rt"); + if(!fileExist(net.c_str())) { + FatalError("The given network does not exist. Create the rt first."); + } + +#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 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())) + 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); + if(n_batch < 1 || n_batch > 64) + FatalError("Batch dim not supported"); + float conf_thresh = YAMLgetConf(conf, "conf_thresh", 0.3); + bool show = YAMLgetConf(conf, "show", true); + bool save = YAMLgetConf(conf, "save", false); + + std::cout <<"Net settings - net: "<< net + <<", ntype: "<< ntype + <<", n_classes: "<< n_classes + <<", n_batch: "<< n_batch + <<", conf_thresh: "<< conf_thresh<<"\n"; + std::cout <<"Demo settings - input: "<< input + <<", show: "<< show + <<", save: "<< save<<"\n\n"; + tk::dnn::Yolo3Detection yolo; tk::dnn::CenternetDetection cnet; tk::dnn::MobilenetDetection mbnet; @@ -98,6 +93,11 @@ int main(int argc, char *argv[]) { FatalError("Network type not allowed (3rd parameter)\n"); } + if(ntype == 'c' || ntype == 'm'){ + cfgPath = ""; + namePath = ""; + } + detNN->init(net,cfgPath,namePath,n_classes,n_batch,conf_thresh); gRun = true; @@ -109,7 +109,7 @@ int main(int argc, char *argv[]) { std::cout<<"camera started\n"; cv::VideoWriter resultVideo; - if(SAVE_RESULT) { + if(save) { int w = cap.get(cv::CAP_PROP_FRAME_WIDTH); int h = cap.get(cv::CAP_PROP_FRAME_HEIGHT); resultVideo.open("result.mp4", cv::VideoWriter::fourcc('M','P','4','V'), 30, cv::Size(w, h)); @@ -149,7 +149,7 @@ int main(int argc, char *argv[]) { cv::waitKey(1); } } - if(n_batch == 1 && SAVE_RESULT) + if(n_batch == 1 && save) resultVideo << frame; } @@ -157,7 +157,7 @@ int main(int argc, char *argv[]) { double mean = 0; std::cout<stats.begin(), detNN->stats.end())/n_batch<<" ms\n"; + std::cout<<"Min: "<<*std::min_element(detNN->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"; for(int i=0; istats.size(); i++) mean += detNN->stats[i]; mean /= detNN->stats.size(); std::cout<<"Avg: "<> /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 -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-9 \ - libtbb-dev \ - libgstreamer1.0-dev \ - libgstreamer-plugins-base1.0-dev \ - libdc1394-22-dev \ - libavresample-dev -RUN cd && wget https://github.com/opencv/opencv/archive/4.5.4.tar.gz && tar -xf 4.5.4.tar.gz && rm *.tar.gz -RUN cd && wget https://github.com/opencv/opencv_contrib/archive/4.5.4.tar.gz && tar -xf 4.5.4.tar.gz && rm *.tar.gz -RUN cd && \ +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.5.4/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/demo.md b/docs/demo.md index b14baf7..bf6b79d 100644 --- a/docs/demo.md +++ b/docs/demo.md @@ -30,31 +30,24 @@ cmake .. -DCMAKE_BUILD_TYPE=Debug -DDEBUG=True make ``` -Once you have successfully created your rt file, run the demo(yolo) : +Once you have successfully created your rt file, run the demo: ``` -./demo yolo4_fp32.rt ../demo/yolo_test.mp4 y 80 ../tests/darknet/cfg/yolo4.cfg ../tests/darknet/names/coco.names +./ 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"```. -To run demo for mobilenet and centernet for the created rt file : -``` -./demo mobilenetv2ssd_fp32.rt m 20 -``` - -In general the demo program takes 7 parameters: -``` -./demo -``` -where - -* `````` is the rt file generated by a test -* ```<``` is the path to a video file or a camera input -* `````` is the type of network. Thee types are currently supported: ```y``` (YOLO family), ```c``` (CenterNet family) and ```m``` (MobileNet-SSD family) -* ``````is the number of classes the network is trained on -* ``` ```is the relative path to the config file (only for darknet based networks) used to train the network -* ``````is the relative path to the names file (only for darknet based networks) used to train the network -* `````` 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). -* `````` if set to 0 the demo will not show the visualization but save the video into result.mp4 (if n-batches ==1) -* `````` confidence threshold for the detector. Only bounding boxes with threshold greater than conf-thresh will be displayed. +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) +* ```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 @@ -69,7 +62,8 @@ To run the demo with FP16 inference follow these steps (example with yolov3): export TKDNN_MODE=FP16 # set the half floating point optimization rm yolo4_fp16.rt # be sure to delete(or move) old tensorRT files ./test_yolo4 # run the yolo test (is slow) -./demo yolo4_fp16.rt ../demo/yolo_test.mp4 y 80 ../tests/darknet/cfg/yolo4.cfg ../tests/darknet/names/coco.names +#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). @@ -94,7 +88,8 @@ export TKDNN_CALIB_LABEL_PATH=../demo/COCO_val2017/all_labels.txt export TKDNN_CALIB_IMG_PATH=../demo/COCO_val2017/all_images.txt rm yolo4_int8.rt # be sure to delete(or move) old tensorRT files ./test_yolo4 # run the yolo test (is slow) -./demo yolo4_int8.rt ../demo/yolo_test.mp4 y 80 ../tests/darknet/cfg/yolo4.cfg ../tests/darknet/names/coco.names +#set net: yolo4_int8.rt in the config file +./demo ``` N.B. diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 5273c83..662eb6a 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 @@ -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"; } } @@ -520,9 +522,33 @@ 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: diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index ddd3d9f..b859074 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -23,6 +23,8 @@ #include #include #include +#include +#include @@ -93,10 +95,16 @@ public: nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Region *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Shortcut *l); nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Yolo *l); - nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Upsample *l); + nvinfer1::IResizeLayer* 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); +#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/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/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/FlattenConcatRT.h b/include/tkDNN/pluginsRT/FlattenConcatRT.h index 02ff596..f7ec495 100644 --- a/include/tkDNN/pluginsRT/FlattenConcatRT.h +++ b/include/tkDNN/pluginsRT/FlattenConcatRT.h @@ -1,3 +1,6 @@ +#ifndef _FLATTENCONCATRT_PLUGIN_H +#define _FLATTENCONCATRT_PLUGIN_H + #include #include #include @@ -93,4 +96,5 @@ namespace nvinfer1 { }; REGISTER_TENSORRT_PLUGIN(FlattenConcatRTPluginCreator); -}; \ No newline at end of file +}; +#endif \ No newline at end of file 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/utils.h b/include/tkDNN/utils.h index 1219ec3..a1a1f1c 100644 --- a/include/tkDNN/utils.h +++ b/include/tkDNN/utils.h @@ -6,6 +6,8 @@ #include #include #include +#include + #include "cuda.h" #include "cuda_runtime_api.h" @@ -16,7 +18,6 @@ #ifdef __linux__ #include - #endif #include @@ -161,5 +162,19 @@ static inline bool isCudaPointer(void *data) { return cudaPointerGetAttributes(&attr, data) == 0; } +inline YAML::Node YAMLloadConf(const std::string& conf_file) { + std::cerr<<"Loading YAML: "< +inline T YAMLgetConf(YAML::Node conf, std::string key, T defaultVal) { + T val = defaultVal; + if(conf && conf[key]) { + val = conf[key].as(); + } + return val; +} + #endif //UTILS_H diff --git a/scripts/checkExecTimes.py b/scripts/checkExecTimes.py new file mode 100644 index 0000000..c3e1b0b --- /dev/null +++ b/scripts/checkExecTimes.py @@ -0,0 +1,37 @@ +import sys +import pandas as pd + +if len(sys.argv) < 3: + print("Error: two csv files are needed, old first new second") + exit(1) + +old_perf_file = str(sys.argv[1]) +new_perf_file = str(sys.argv[2]) + +verbose = False +if len(sys.argv) == 4: + verbose = bool(sys.argv[3]) + +print("Comparing {} vs {}".format(old_perf_file, new_perf_file)) + +df_old = pd.read_csv (old_perf_file, sep=';', header=None, index_col=0) +df_new = pd.read_csv (new_perf_file, sep=';', header=None, index_col=0) + +for index, row in df_new.iterrows(): + if index in df_old.index: + if verbose: + print("New: ",row[1], row[2], row[3]) + print("Old: ",df_old.loc[index][1], df_old.loc[index][2], df_old.loc[index][3]) + + print(index, end=': ') + if abs(row[1] - df_old.loc[index][1]) < df_old.loc[index][1]*0.1: + print("similar performance") + elif (row[1] < df_old.loc[index][1]): + print('\x1b[3;30;42m' + 'faster' + '\x1b[0m') + elif (row[1] > df_old.loc[index][1]): + if row[1] > df_old.loc[index][1] + df_old.loc[index][1] * 0.5 : + print('\x1b[3;30;41m' + 'WAY SLOWER' + '\x1b[0m') + else: + print('\x1b[3;30;41m' + 'slower' + '\x1b[0m') + + diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 2e9eb6e..3f086bd 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -26,15 +26,15 @@ class Logger : public ILogger { namespace tk { namespace dnn { -std::maptensors; +std::maptensors; NetworkRT::NetworkRT(Network *net, const char *name) { - float rt_ver = float(NV_TENSORRT_MAJOR) + - float(NV_TENSORRT_MINOR)/10 + + float rt_ver = float(NV_TENSORRT_MAJOR) + + float(NV_TENSORRT_MINOR)/10 + float(NV_TENSORRT_PATCH)/100; std::cout<<"New NetworkRT (TensorRT v"<platformHasFastFp16()<<"\n"; std::cout<<"Int8 support: "<platformHasFastInt8()<<"\n"; @@ -42,12 +42,12 @@ NetworkRT::NetworkRT(Network *net, const char *name) { std::cout<<"DLAs: "<getNbDLACores()<<"\n"; #endif networkRT = builderRT->createNetworkV2(0U); -#if NV_TENSORRT_MAJOR >= 6 +#if NV_TENSORRT_MAJOR >= 6 configRT = builderRT->createBuilderConfig(); #endif - + if(!fileExist(name)) { -#if NV_TENSORRT_MAJOR >= 6 +#if NV_TENSORRT_MAJOR >= 6 // Calibrator life time needs to last until after the engine is built. std::unique_ptr calibrator; @@ -78,14 +78,14 @@ NetworkRT::NetworkRT(Network *net, const char *name) { configRT->setDLACore(0); } #endif -#if NV_TENSORRT_MAJOR >= 6 +#if NV_TENSORRT_MAJOR >= 6 if(net->int8 && builderRT->platformHasFastInt8()){ // dtRT = DataType::kINT8; // builderRT->setInt8Mode(true); configRT->setFlag(BuilderFlag::kINT8); - BatchStream calibrationStream(dim, 1, 100, //TODO: check if 100 images are sufficient to the calibration (or 4951) + BatchStream calibrationStream(dim, 1, 100, //TODO: check if 100 images are sufficient to the calibration (or 4951) net->fileImgList, net->fileLabelList); - + /* The calibTableFilePath contains the path+filename of the calibration table. * Each calibration table can be found in the corresponding network folder (../Test/*). * Each network is located in a folder with the same name as the network. @@ -96,15 +96,15 @@ NetworkRT::NetworkRT(Network *net, const char *name) { if(!fileExist((const char *)calib_table_path.c_str())) calib_table_name = "./" + net->networkNameRT.substr(0, net->networkNameRT.find('.')) + "-calibration.table"; - calibrator.reset(new Int8EntropyCalibrator(calibrationStream, 1, - calib_table_name, + calibrator.reset(new Int8EntropyCalibrator(calibrationStream, 1, + calib_table_name, "data")); configRT->setInt8Calibrator(calibrator.get()); } #endif - + // add input layer - ITensor *input = networkRT->addInput("data", DataType::kFLOAT, + ITensor *input = networkRT->addInput("data", DataType::kFLOAT, Dims3{ dim.c, dim.h, dim.w}); checkNULL(input); @@ -112,17 +112,17 @@ NetworkRT::NetworkRT(Network *net, const char *name) { for(int i=0; inum_layers; i++) { Layer *l = net->layers[i]; ILayer *Ilay = convert_layer(input, l); -#if NV_TENSORRT_MAJOR >= 6 +#if NV_TENSORRT_MAJOR >= 6 if(net->int8 && builderRT->platformHasFastInt8()) { Ilay->setPrecision(DataType::kINT8); } #endif Ilay->setName( (l->getLayerName() + std::to_string(i)).c_str() ); - + input = Ilay->getOutput(0); input->setName( (l->getLayerName() + std::to_string(i) + "_out").c_str() ); - + if(l->final) networkRT->markOutput(*input); tensors[l] = input; @@ -137,12 +137,16 @@ NetworkRT::NetworkRT(Network *net, const char *name) { std::cout<<"Selected maxBatchSize: "<getMaxBatchSize()<<"\n"; printCudaMemUsage(); std::cout<<"Building tensorRT cuda engine...\n"; -#if NV_TENSORRT_MAJOR >= 6 +#if NV_TENSORRT_MAJOR >= 6 && NV_TENSORRT_MAJOR <=7 engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT); -#else +#elif NV_TENSORRT_MAJOR < 6 engineRT = builderRT->buildCudaEngine(*networkRT); //engineRT = std::shared_ptr(builderRT->buildCudaEngine(*networkRT)); +#elif NV_TENSORRT_MAJOR >=8 + IHostMemory *serializedEngineRT = builderRT->buildSerializedNetwork(*networkRT,*configRT); + #endif +#if NV_TENSORRT_MAJOR > 5 && NV_TENSORRT_MAJOR < 8 if(engineRT == nullptr) FatalError("cloud not build cuda engine") // we don't need the network any more @@ -150,6 +154,19 @@ NetworkRT::NetworkRT(Network *net, const char *name) { std::cout<<"serialize net\n"; builderActive = true; serialize(name); +#else + if(serializedEngineRT == nullptr){ + FatalError("could not build cuda engine"); + } + std::cout<<"saving serialized network to file"<= 8 + deserialize(name); +#endif + +#endif } else { builderActive = false; deserialize(name); @@ -165,7 +182,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) { // In order to bind the buffers, we need to know the names of the input and output tensors. // note that indices are guaranteed to be less than IEngine::getNbBindings() - buf_input_idx = engineRT->getBindingIndex("data"); + buf_input_idx = engineRT->getBindingIndex("data"); buf_output_idx = engineRT->getBindingIndex("out"); std::cout<<"input index = "< output index = "<getLayerName()<<"\n"; FatalError("Layer not implemented in tensorRT"); @@ -268,10 +287,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Dense *l) { //std::cout<<"convert Dense\n"; void *data_b, *bias_b; if(dtRT == DataType::kHALF) { - data_b = l->data16_h; + data_b = l->data16_h; bias_b = l->bias16_h; } else { - data_b = l->data_h; + data_b = l->data_h; bias_b = l->bias_h; } @@ -291,7 +310,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { void *data_b, *bias_b, *bias2_b, *power_b, *mean_b, *variance_b, *scales_b; if(dtRT == DataType::kHALF) { - data_b = l->data16_h; + data_b = l->data16_h; bias_b = l->bias16_h; bias2_b = l->bias216_h; power_b = l->power16_h; @@ -299,7 +318,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { variance_b = l->variance16_h; scales_b = l->scales16_h; } else { - data_b = l->data_h; + data_b = l->data_h; bias_b = l->bias_h; bias2_b = l->bias2_h; power_b = l->power_h; @@ -315,14 +334,15 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { b = { dtRT, bias_b, l->outputs}; else{ if (l->additional_bias) - b = { dtRT, bias2_b, l->outputs}; + b = { dtRT, bias2_b, l->outputs}; else b = { dtRT, nullptr, 0}; //on batchnorm bias are added later } ILayer *lRT = nullptr; +#if NV_TENSORRT_MAJOR < 8 if(!l->deConv) { - IConvolutionLayer *lRTconv = networkRT->addConvolution(*input, + IConvolutionLayer *lRTconv = networkRT->addConvolution(*input, l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b); checkNULL(lRTconv); lRTconv->setStride(DimsHW{l->strideH, l->strideW}); @@ -330,17 +350,39 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { lRTconv->setNbGroups(l->groups); lRT = (ILayer*) lRTconv; } else { - IDeconvolutionLayer *lRTconv = networkRT->addDeconvolution(*input, + IDeconvolutionLayer *lRTconv = networkRT->addDeconvolution(*input, l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b); checkNULL(lRTconv); lRTconv->setStride(DimsHW{l->strideH, l->strideW}); lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW}); lRTconv->setNbGroups(l->groups); lRT = (ILayer*) lRTconv; - + Dims d = lRTconv->getOutput(0)->getDimensions(); //std::cout<<"DECONV: "<deConv) { + IConvolutionLayer *lRTconv = networkRT->addConvolutionNd(*input, + l->outputs, Dims2{l->kernelH, l->kernelW}, w, b); + checkNULL(lRTconv); + lRTconv->setStrideNd(Dims2{l->strideH, l->strideW}); + lRTconv->setPaddingNd(Dims2{l->paddingH, l->paddingW}); + lRTconv->setNbGroups(l->groups); + lRT = (ILayer*) lRTconv; + } else { + IDeconvolutionLayer *lRTconv = networkRT->addDeconvolutionNd(*input, + l->outputs, Dims2{l->kernelH, l->kernelW}, w, b); + checkNULL(lRTconv); + lRTconv->setStrideNd(Dims2{l->strideH, l->strideW}); + lRTconv->setPaddingNd(Dims2{l->paddingH, l->paddingW}); + lRTconv->setNbGroups(l->groups); + lRT = (ILayer*) lRTconv; + + Dims d = lRTconv->getOutput(0)->getDimensions(); + //std::cout<<"DECONV: "<batchnorm) { @@ -348,14 +390,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { Weights shift{dtRT, mean_b, l->outputs}; Weights scale{dtRT, variance_b, l->outputs}; // std::cout<getNbOutputs()<addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, + IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, shift, scale, power); - + checkNULL(lRT2); Weights shift2{dtRT, bias_b, l->outputs}; Weights scale2{dtRT, scales_b, l->outputs}; - IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, + IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, shift2, scale2, power); checkNULL(lRT3); @@ -396,13 +438,83 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) { } else { +#if NV_TENSORRT_MAJOR < 8 IPoolingLayer *lRT = networkRT->addPooling(*input, ptype, DimsHW{l->winH, l->winW}); checkNULL(lRT); lRT->setPadding(DimsHW{l->paddingH, l->paddingW}); lRT->setStride(DimsHW{l->strideH, l->strideW}); return lRT; - } +#else + IPoolingLayer *lRT = networkRT->addPoolingNd(*input,ptype,Dims2{l->winH,l->winW}); + checkNULL(lRT); + lRT->setPaddingNd(Dims2{l->paddingH,l->paddingW}); + lRT->setStrideNd(Dims2{l->strideH,l->strideW}); + return lRT; +#endif + } +} + +ILayer* NetworkRT::convert_layer(ITensor *input,Padding *l){ + + float rt_ver = float(NV_TENSORRT_MAJOR) + + float(NV_TENSORRT_MINOR)/10 + + float(NV_TENSORRT_PATCH)/100; + +#if ((NV_TENSORRT_MAJOR == 8 && NV_TENSORRT_MINOR >= 2) || NV_TENSORRT_MAJOR > 8) + auto *lRT = networkRT->addSlice(*input,Dims3{0,0,0},Dims3{l->output_dim.c,l->output_dim.h,l->output_dim.w},Dims3{0,0,0}); + if(l->padding_mode == PADDING_MODE_REFLECTION){ + lRT->setMode(SliceMode::kREFLECT); + }else if(l->padding_mode == PADDING_MODE_CONSTANT || l->padding_mode == PADDING_MODE_ZERO){ + lRT->setMode(SliceMode::kFILL); + lRT->setInput(4, reinterpret_cast(l->constant)); + } + checkNULL(lRT); + return lRT; +#else + //todo add PADDING_MODE_CONSTANT AND PADDING_MODE_ZERO for tensorrt versions < 8.2 + if(l->padding_mode == PADDING_MODE_REFLECTION){ + auto creator = getPluginRegistry()->getPluginCreator("ReflectionPaddingRT_tkDNN","1"); + std::vector mPluginAttributes; + PluginFieldCollection mFC{}; + mPluginAttributes.emplace_back(PluginField("padH",&l->paddingH,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("padW",&l->paddingW,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("inputH",&l->input_dim.h,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("inputW",&l->input_dim.w,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("outputH",&l->output_dim.h,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("outputW",&l->output_dim.w,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("n",&l->input_dim.n,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("c",&l->input_dim.c,PluginFieldType::kINT32,1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); + auto *plugin = creator->createPlugin(l->getLayerName().c_str(),&mFC); + auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); + checkNULL(lRT); + return lRT; + }else if(l->padding_mode == PADDING_MODE_CONSTANT || l->padding_mode == PADDING_MODE_ZERO){ + auto creator = getPluginRegistry()->getPluginCreator("ConstantPaddingRT_tkDNN","1"); + std::vector mPluginAttributes; + PluginFieldCollection mFC{}; + mPluginAttributes.emplace_back(PluginField("padH",&l->paddingH,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("padW",&l->paddingW,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("inputH",&l->input_dim.h,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("inputW",&l->input_dim.w,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("outputH",&l->output_dim.h,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("outputW",&l->output_dim.w,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("n",&l->input_dim.n,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("c",&l->input_dim.c,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("constant",&l->constant,PluginFieldType::kFLOAT32,1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); + auto *plugin = creator->createPlugin(l->getLayerName().c_str(),&mFC); + auto *lRT = networkRT->addPluginV2(&input,1,*plugin); + checkNULL(lRT); + return lRT; + } + + return nullptr; + +#endif } ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { @@ -410,14 +522,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { if(l->act_mode == ACTIVATION_LEAKY) { //std::cout<<"New plugin LEAKY\n"; - -#if NV_TENSORRT_MAJOR < 6 + +#if NV_TENSORRT_MAJOR < 6 // plugin version IPlugin *plugin = new ActivationLeakyRT(l->slope); IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); checkNULL(lRT); return lRT; -#else +#else IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kLEAKY_RELU); lRT->setAlpha(l->slope); checkNULL(lRT); @@ -442,7 +554,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { //IPluginV2Layer *lRT = networkRT->addPluginV2(&input, 1, *plugin); //checkNULL(lRT); return lRT; - } + } else if(l->act_mode == ACTIVATION_MISH) { IActivationLayer *lRT1 = networkRT->addActivation(*input, ActivationType::kSOFTPLUS); lRT1->setAlpha(1); @@ -456,6 +568,11 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { checkNULL(lRT); return lRT; } + else if(l->act_mode == CUDNN_ACTIVATION_ELU || l->act_mode == ACTIVATION_ELU){ + IActivationLayer *lRT = networkRT->addActivation(*input,ActivationType::kELU); + checkNULL(lRT); + return lRT; + } else { FatalError("this Activation mode is not yet implemented"); return NULL; @@ -473,7 +590,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Softmax *l) { ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) { // std::cout<<"convert route\n"; - + ITensor **tens = new ITensor*[l->layers_n]; @@ -585,10 +702,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) { //std::cout<<"convert Shortcut\n"; //std::cout<<"New plugin Shortcut\n"; - + ITensor *back_tens = tensors[l->backLayer]; - if(l->backLayer->output_dim.c == l->output_dim.c && !l->mul) + if(l->backLayer->output_dim.c == l->output_dim.c && !l->mul) { IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM); checkNULL(lRT); @@ -612,7 +729,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) { auto *plugin = creator->createPlugin(l->getLayerName().c_str(),&mFC); auto **inputs = new ITensor*[2]; inputs[0] = input; - inputs[1] = back_tens; + inputs[1] = back_tens; auto *lRT = networkRT->addPluginV2(inputs, 2, *plugin); checkNULL(lRT); return lRT; @@ -642,8 +759,9 @@ IPluginV2Layer* NetworkRT::convert_layer(ITensor *input, Yolo *l) { return lRT; } -IPluginV2Layer* NetworkRT::convert_layer(ITensor *input, Upsample *l) { - //std::cout<<"convert Upsample\n"; +IResizeLayer* NetworkRT::convert_layer(ITensor *input, Upsample *l) { + +#if NV_TENSORRT_MAJOR < 8 auto creator = getPluginRegistry()->getPluginCreator("UpSample_tkDNN","1"); std::vector mPluginAttributes; PluginFieldCollection mFC{}; @@ -657,6 +775,13 @@ IPluginV2Layer* NetworkRT::convert_layer(ITensor *input, Upsample *l) { auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; +#else + auto *lRT = networkRT->addResize(*input); + lRT->setResizeMode(ResizeMode::kNEAREST); + lRT->setOutputDimensions(Dims3{l->output_dim.c, l->output_dim.h, l->output_dim.w}); + checkNULL(lRT); + return lRT; +#endif } ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) { @@ -739,20 +864,21 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) { Weights shift{dtRT, mean_b, l->outputs}; Weights scale{dtRT, variance_b, l->outputs}; //std::cout<getNbOutputs()<addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, + IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, shift, scale, power); - + checkNULL(lRT2); Weights shift2{dtRT, bias_b, l->outputs}; Weights scale2{dtRT, scales_b, l->outputs}; - IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, + IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, shift2, scale2, power); checkNULL(lRT3); return lRT3; } +#if NV_TENSORRT_MAJOR > 5 && NV_TENSORRT_MAJOR < 8 bool NetworkRT::serialize(const char *filename) { std::ofstream p(filename, std::ios::binary); @@ -769,6 +895,21 @@ bool NetworkRT::serialize(const char *filename) { ptr->destroy(); return true; } +#else +bool NetworkRT::serialize(const char *filename,nvinfer1::IHostMemory *ptr){ + std::ofstream p(filename, std::ios::binary); + if (!p) { + FatalError("could not open plan output file"); + return false; + } + + if(ptr == nullptr) + FatalError("Cant serialize network"); + + p.write(reinterpret_cast(ptr->data()), ptr->size()); + return true; +} +#endif bool NetworkRT::deserialize(const char *filename) { @@ -793,10 +934,10 @@ bool NetworkRT::deserialize(const char *filename) { } void NetworkRT::destroy() { - contextRT->destroy(); + delete contextRT; if(builderActive) { - engineRT->destroy(); - builderRT->destroy(); + delete engineRT; + delete builderRT; } } diff --git a/src/Padding.cpp b/src/Padding.cpp new file mode 100644 index 0000000..1ba6d38 --- /dev/null +++ b/src/Padding.cpp @@ -0,0 +1,45 @@ +// +// Created by perseusdg on 03/01/22. +// + +#include +#include "Layer.h" +#include "kernels.h" + +namespace tk{ namespace dnn { + Padding::Padding(Network *net, int32_t pad_h, int32_t pad_w, tkdnnPaddingMode_t padding_mode,float constant) : Layer(net) { + this->paddingH = pad_h; + this->paddingW = pad_w; + this->padding_mode = padding_mode; + output_dim.c = input_dim.c; + output_dim.n = input_dim.n; + output_dim.h = input_dim.h + 2 * (this->paddingH); + output_dim.w = input_dim.w + 2 * (this->paddingW); + if(padding_mode == tkdnnPaddingMode_t::PADDING_MODE_CONSTANT){ + this->constant = constant; + }else{ + this->constant = 0; + } + checkCuda(cudaMalloc(&dstData,output_dim.tot()*sizeof(dnnType))); + } + + Padding::~Padding() { + checkCuda(cudaFree(dstData)); + } + dnnType* Padding::infer(dataDim_t &dim, float *srcData) { + fill(dstData,output_dim.tot(),0.0); + if(padding_mode == tkdnnPaddingMode_t::PADDING_MODE_REFLECTION) + { + reflection_pad2d_out_forward(paddingH, paddingW, srcData, dstData, input_dim.h, input_dim.w, input_dim.c, + input_dim.n); + } + else if(padding_mode == tkdnnPaddingMode_t::PADDING_MODE_CONSTANT){ + constant_pad2d_forward(srcData,dstData,input_dim.h,input_dim.w,output_dim.h,output_dim.w,input_dim.c, + input_dim.n,paddingH,paddingW,constant); + } + + dim = output_dim; + return dstData; + } + +}} diff --git a/src/kernels/padding.cu b/src/kernels/padding.cu new file mode 100644 index 0000000..eafcd5e --- /dev/null +++ b/src/kernels/padding.cu @@ -0,0 +1,110 @@ +#include "kernels.h" +#include +#include + +/* + * Reflection padding is from https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/native/cuda/ReflectionPad.cu + */ +__device__ +inline thrust::pair get_index_mapping2d( + int32_t input_dim_x,int32_t input_dim_y,int32_t output_dim_x, + int32_t output_dim_y,int32_t pad_l,int32_t pad_t,int32_t output_xy, + int32_t y_shift,int32_t z_shift,int32_t n_plane){ + auto input_offset = ((blockIdx.y + y_shift) + (blockIdx.z + z_shift)*n_plane)*input_dim_x*input_dim_y; + auto output_offset = ((blockIdx.y + y_shift) + (blockIdx.z + z_shift)*n_plane)*output_dim_x*output_dim_y; + auto output_x = output_xy % output_dim_x; + auto output_y = output_xy/output_dim_x; + + auto i_start_x = ::max(int32_t(0),-pad_l); + auto i_start_y = ::max(int32_t(0),-pad_t); + auto o_start_x = ::max(int32_t(0),pad_l); + auto o_start_y = ::max(int32_t(0),pad_t); + + auto input_x = ::abs(output_x - pad_l) - ::abs(output_x - (input_dim_x + pad_l -1)) -output_x + 2*pad_l + input_dim_x -1 -o_start_x + i_start_x; + auto input_y = ::abs(output_y - pad_t) - ::abs(output_y - (input_dim_y + pad_t -1)) -output_y + 2*pad_t + input_dim_y -1 -o_start_y + i_start_y; + + return thrust::make_pair(input_offset + input_y*input_dim_x + input_x,output_offset + output_y*output_dim_x+output_x); +} + +__global__ +void reflection_pad2d_out_kernel( + float* input,float* output,int32_t input_dim_x, + int32_t input_dim_y,int32_t pad_t,int32_t pad_b,int32_t pad_l, + int32_t pad_r,int32_t y_shift,int32_t z_shift,int32_t n_plane){ + auto output_xy = threadIdx.x + blockIdx.x * blockDim.x; + auto output_dim_x = input_dim_x + pad_l + pad_r; + auto output_dim_y = input_dim_y + pad_t + pad_b; + + if(output_xy < output_dim_x*output_dim_y){ + auto index_pair = get_index_mapping2d(input_dim_x,input_dim_y,output_dim_x,output_dim_y,pad_l,pad_t,output_xy,y_shift,z_shift,n_plane); + output[index_pair.second] = input[index_pair.first]; + } +} + +int32_t ceilDiv(int32_t a,int32_t b){ + return (a+b-1)/b; +} + + +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){ + int32_t pad_l = pad_w; + int32_t pad_r = pad_w; + int32_t pad_t = pad_h; + int32_t pad_b = pad_w; + int32_t output_h = input_h + pad_t + pad_b; + int32_t output_w = input_w + pad_l + pad_r; + int32_t size_y = plane_dim; + int32_t size_z = n_batch; + int32_t output_plane_size = output_h*output_w; + dim3 block_size(output_plane_size>256 ?256:output_plane_size); + for(int32_t block_y=0;block_y(65535)); + for(int32_t block_z=0;block_z(65535)); + + dim3 grid_size(ceilDiv(output_plane_size,static_cast(256)),block_y_size,block_z_size); + reflection_pad2d_out_kernel<<>>(srcData,dstData,input_w,input_h,pad_t,pad_b,pad_l,pad_r,block_y,block_z,plane_dim); + } + } + +} + +/* + * constant padding is inspired from https://github.com/apache/incubator-mxnet/blob/master/src/operator/pad.cu + */ + +__global__ +void constant_pad2d_kernel(dnnType *srcData,dnnType *dstData,const int32_t padT,const int32_t padL,float constant,int32_t n,int32_t c,int32_t i_h,int32_t i_w,int32_t o_h,int32_t o_w){ + int outputPointId = threadIdx.x + blockIdx.x * blockDim.x; + if(outputPointId >= o_h*o_w){ + return ; + } + + int Ny = i_h; + int Nx = i_w; + + int plane = blockIdx.y; + int batch = blockIdx.z; + int outputPointX = outputPointId % o_w; + int outputPointY = outputPointId / o_w; + int checkT = max(0, outputPointY - padT + 1); + int checkB = max(0, padT + Ny - outputPointY); + int checkL = max(0, outputPointX - padL + 1); + int checkR = max(0, padL + Nx - outputPointX); + int inputPointX = min(max(outputPointX - padL, 0), Nx - 1); + int inputPointY = min(max(outputPointY - padT, 0), Ny - 1); + int need_pad = !(checkT * checkB * checkL * checkR); + float value_to_copy = srcData[batch*c*i_h*i_w + plane*i_h*i_w + inputPointY*i_w + inputPointX]; + dstData[batch*c*o_w*o_h + plane*o_h*o_w + outputPointY*o_w + outputPointX] = value_to_copy * (!need_pad) + need_pad*constant; + +} + +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){ + int32_t output_plane_size = output_h*output_w; + dim3 block_size(output_plane_size>256 ?256:output_plane_size); + dim3 grid_size(ceilDiv(output_plane_size,static_cast(256)),c,n); + constant_pad2d_kernel<<>>(srcData,dstData,padT,padL,constant,n,c,input_h,input_w,output_h,output_w); + +} + diff --git a/src/pluginsRT/ConstantPaddingRT.cpp b/src/pluginsRT/ConstantPaddingRT.cpp new file mode 100644 index 0000000..ac37e9d --- /dev/null +++ b/src/pluginsRT/ConstantPaddingRT.cpp @@ -0,0 +1,201 @@ +#include + +using namespace nvinfer1; + +std::vector ConstantPaddingRTPluginCreator::mPluginAttributes; +PluginFieldCollection ConstantPaddingRTPluginCreator::mFC{}; + +static const char* CONSTANTPADDINGRT_PLUGIN_VERSION{"1"}; +static const char* CONSTANTPADDINGRT_PLUGIN_NAME{"ConstantPaddingRT_tkDNN"}; + +ConstantPaddingRT::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) { + this->padH = padH; + this->padW = padW; + this->n = n; + this->c = c; + this->i_h = i_h; + this->i_w = i_w; + this->o_h = o_h; + this->o_w = o_w; + this->constant = constant; + +} + +ConstantPaddingRT::ConstantPaddingRT(const void *data, size_t length) { + const char* buf = reinterpret_cast(data),*bufcheck=buf; + padH = readBUF(buf); + padW = readBUF(buf); + i_h = readBUF(buf); + i_w = readBUF(buf); + o_h = readBUF(buf); + o_w = readBUF(buf); + n = readBUF(buf); + c = readBUF(buf); + constant = readBUF(buf); + assert(buf = bufcheck + length); +} + +ConstantPaddingRT::~ConstantPaddingRT() {} + +int ConstantPaddingRT::getNbOutputs() const NOEXCEPT{ + return 1; +} + +Dims ConstantPaddingRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return Dims3{c,o_h,o_w}; +} + +int ConstantPaddingRT::initialize() NOEXCEPT { + return 0; +} + +void ConstantPaddingRT::terminate() NOEXCEPT { + +} + +size_t ConstantPaddingRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { + return 0; +} + +#if NV_TENSORRT_MAJOR > 7 +int ConstantPaddingRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT { + dnnType* srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType* dstData = reinterpret_cast(outputs[0]); + constant_pad2d_forward(srcData,dstData,i_h,i_w,o_h,o_w,c,n,padH,padW,constant,stream); + return 0; +} +#elif NV_TENSORRT_MAJOR <= 7 + int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) { + dnnType* srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType* dstData = reinterpret_cast(outputs[0]); + constant_pad2d_forward(srcData,dstData,i_h,i_w,o_h,o_w,c,n,padH,padW,constant,stream); + return 0; +} +#endif + +size_t ConstantPaddingRT::getSerializationSize() const NOEXCEPT { + return (8*sizeof(int32_t) + 1*sizeof(float)); +} + +void ConstantPaddingRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer),*a=buf; + writeBUF(buf,padH); + writeBUF(buf,padW); + writeBUF(buf,i_h); + writeBUF(buf,i_w); + writeBUF(buf,o_h); + writeBUF(buf,o_w); + writeBUF(buf,n); + writeBUF(buf,c); + writeBUF(buf,constant); +} + +void ConstantPaddingRT::destroy() NOEXCEPT { + delete this; +} + +const char* ConstantPaddingRT::getPluginType() const NOEXCEPT { + return CONSTANTPADDINGRT_PLUGIN_NAME; +} + +const char* ConstantPaddingRT::getPluginVersion() const NOEXCEPT { + return CONSTANTPADDINGRT_PLUGIN_VERSION; +} + +const char* ConstantPaddingRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void ConstantPaddingRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +IPluginV2Ext *ConstantPaddingRT::clone() const NOEXCEPT { + auto *p = new ConstantPaddingRT(padH,padW,n,c,i_h,i_w,o_h,o_w,constant); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + +DataType ConstantPaddingRT::getOutputDataType(int index, const nvinfer1::DataType *inputTypes, + int nbInputs) const NOEXCEPT { + return DataType::kFLOAT; +} + +void ConstantPaddingRT::attachToContext(cudnnContext *cudnnContext, cublasContext *cublasContext, + IGpuAllocator *gpuAllocator) NOEXCEPT { + +} + +bool ConstantPaddingRT::isOutputBroadcastAcrossBatch(int outputIndex, const bool *inputIsBroadcasted, + int nbInputs) const NOEXCEPT { + return false; +} + +bool ConstantPaddingRT::canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT { + return false; +} + +void ConstantPaddingRT::configurePlugin(const Dims *inputDims, int32_t nbInputs, const Dims *outputDims, + int32_t nbOutputs, const DataType *inputTypes, const DataType *outputTypes, + const bool *inputIsBroadcast, const bool *outputIsBroadcast, + PluginFormat floatFormat, int32_t maxBatchSize) NOEXCEPT { + +} + +void ConstantPaddingRT::detachFromContext() NOEXCEPT { + +} + +bool ConstantPaddingRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR); +} + +ConstantPaddingRTPluginCreator::ConstantPaddingRTPluginCreator() { + mPluginAttributes.clear(); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void ConstantPaddingRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *ConstantPaddingRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2Ext *ConstantPaddingRTPluginCreator::deserializePlugin(const char *name, const void *serialData, + size_t serialLength) NOEXCEPT { + auto *pluginObj = new ConstantPaddingRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2Ext *ConstantPaddingRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + int padH = *(static_cast(fields[0].data)); + int padW = *(static_cast(fields[1].data)); + int inputH = *(static_cast(fields[2].data)); + int inputW = *(static_cast(fields[3].data)); + int outputH = *(static_cast(fields[4].data)); + int outputW = *(static_cast(fields[5].data)); + int n = *(static_cast(fields[6].data)); + int c = *(static_cast(fields[7].data)); + float constant = *(static_cast(fields[8].data)); + auto *pluginObj = new ConstantPaddingRT(padH,padW,n,c,inputH,inputW,outputH,outputW,constant); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *ConstantPaddingRTPluginCreator::getPluginName() const NOEXCEPT { + return CONSTANTPADDINGRT_PLUGIN_NAME; +} + +const char *ConstantPaddingRTPluginCreator::getPluginVersion() const NOEXCEPT { + return CONSTANTPADDINGRT_PLUGIN_VERSION; +} + +const PluginFieldCollection *ConstantPaddingRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} diff --git a/src/pluginsRT/ReflectionPadding.cpp b/src/pluginsRT/ReflectionPadding.cpp new file mode 100644 index 0000000..21d897a --- /dev/null +++ b/src/pluginsRT/ReflectionPadding.cpp @@ -0,0 +1,202 @@ +#include +using namespace nvinfer1; + +std::vector ReflectionPaddingRTPluginCreator::mPluginAttributes; +PluginFieldCollection ReflectionPaddingRTPluginCreator::mFC{}; + +static const char* REFLECTIONPADDINGRT_PLUGIN_VERSION{"1"}; +static const char* REFLECTIONPADDINGRT_PLUGIN_NAME{"ReflectionPaddingRT_tkDNN"}; + +ReflectionPaddingRT::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) { + this->padH = padH; + this->padW = padW; + this->input_h = input_h; + this->input_w = input_w; + this->output_h = output_h; + this->output_w = output_w; + this->n = n; + this->c = c; +} + +ReflectionPaddingRT::ReflectionPaddingRT(const void *data, size_t length) { + const char* buf = reinterpret_cast(data),*bufcheck=buf; + padH = readBUF(buf); + padW = readBUF(buf); + input_h = readBUF(buf); + input_w = readBUF(buf); + output_h = readBUF(buf); + output_w = readBUF(buf); + n = readBUF(buf); + c = readBUF(buf); + assert(buf = bufcheck + length); +} + +ReflectionPaddingRT::~ReflectionPaddingRT() {} + +int ReflectionPaddingRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims ReflectionPaddingRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return Dims3{c,output_h,output_w}; +} + +int ReflectionPaddingRT::initialize() NOEXCEPT { + return 0; +} + +void ReflectionPaddingRT::terminate() NOEXCEPT { + +} + +size_t ReflectionPaddingRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { + return 0; +} + +#if NV_TENSORRT_MAJOR > 7 +int ReflectionPaddingRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + dnnType* srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType* dstData = reinterpret_cast(outputs[0]); + reflection_pad2d_out_forward(padH,padW,srcData,dstData,input_h,input_w,c,n,stream); + return 0; +} + +#elif NV_TENSORRT_MAJOR <= 7 +int32_t ReflectionPaddingRT::enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream){ + dnnType* srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType* dstData = reinterpret_cast(outputs[0]); + reflection_pad2d_out_forward(padH,padW,srcData,dstData,input_h,input_w,c,n,stream); + return 0; +} +#endif + + +size_t ReflectionPaddingRT::getSerializationSize() const NOEXCEPT { + return 8*sizeof(int32_t); +} + +void ReflectionPaddingRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer),*a=buf; + writeBUF(buf,padH); + writeBUF(buf,padW); + writeBUF(buf,input_h); + writeBUF(buf,input_w); + writeBUF(buf,output_h); + writeBUF(buf,output_w); + writeBUF(buf,n); + writeBUF(buf,c); +} + +void ReflectionPaddingRT::destroy() NOEXCEPT { + delete this; +} + +const char *ReflectionPaddingRT::getPluginType() const NOEXCEPT { + return REFLECTIONPADDINGRT_PLUGIN_NAME; +} + +const char *ReflectionPaddingRT::getPluginVersion() const NOEXCEPT { + return REFLECTIONPADDINGRT_PLUGIN_VERSION; +} + +const char *ReflectionPaddingRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void ReflectionPaddingRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +IPluginV2Ext *ReflectionPaddingRT::clone() const NOEXCEPT { + auto *p = new ReflectionPaddingRT(padH,padW,input_h,input_w,output_h,output_w,c,n); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + +DataType +ReflectionPaddingRT::getOutputDataType(int index, const nvinfer1::DataType *inputTypes, int nbInputs) const NOEXCEPT { + return DataType::kFLOAT; +} + +void ReflectionPaddingRT::attachToContext(cudnnContext *cudnnContext, cublasContext *cublasContext, + IGpuAllocator *gpuAllocator) NOEXCEPT { +} + +bool ReflectionPaddingRT::isOutputBroadcastAcrossBatch(int outputIndex, const bool *inputIsBroadcasted, + int nbInputs) const NOEXCEPT { + return false; +} + +bool ReflectionPaddingRT::canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT { + return false; +} + +void +ReflectionPaddingRT::configurePlugin(const Dims *inputDims, int32_t nbInputs, const Dims *outputDims, int32_t nbOutputs, + const DataType *inputTypes, const DataType *outputTypes, + const bool *inputIsBroadcast, const bool *outputIsBroadcast, + PluginFormat floatFormat, int32_t maxBatchSize) NOEXCEPT { + +} + +void ReflectionPaddingRT::detachFromContext() NOEXCEPT { + +} + +bool ReflectionPaddingRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR); +} + + +ReflectionPaddingRTPluginCreator::ReflectionPaddingRTPluginCreator() { + mPluginAttributes.clear(); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void ReflectionPaddingRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *ReflectionPaddingRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2Ext *ReflectionPaddingRTPluginCreator::deserializePlugin(const char *name, const void *serialData, + size_t serialLength) NOEXCEPT { + auto *pluginObj = new ReflectionPaddingRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2Ext * +ReflectionPaddingRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + int padH = *(static_cast(fields[0].data)); + int padW = *(static_cast(fields[1].data)); + int inputH = *(static_cast(fields[2].data)); + int inputW = *(static_cast(fields[3].data)); + int outputH = *(static_cast(fields[4].data)); + int outputW = *(static_cast(fields[5].data)); + int n = *(static_cast(fields[6].data)); + int c = *(static_cast(fields[7].data)); + auto *pluginObj = new ReflectionPaddingRT(padH,padW,inputH,inputW,outputH,outputW,c,n); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *ReflectionPaddingRTPluginCreator::getPluginName() const NOEXCEPT { + return REFLECTIONPADDINGRT_PLUGIN_NAME; +} + +const char *ReflectionPaddingRTPluginCreator::getPluginVersion() const NOEXCEPT { + return REFLECTIONPADDINGRT_PLUGIN_VERSION; +} + +const PluginFieldCollection *ReflectionPaddingRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + +