mmr/tkDNN
Archived
1
0
Fork 0
Code Issues 46 Pull Requests 12 Releases 4 Activity

Merge pull request #3 from perseusdg/tensorrt8

Browse Source 
push tensorrt8 commits to the rds branch
This commit is contained in:
Harshvardhan Chandirasekar
2022-01-20 21:12:49 +05:30
committed by GitHub
parent 19f12d9c15 3023926695
commit 53725c88a9
24 changed files with 1442 additions and 181 deletions
Download Patch File Download Diff File Expand all files Collapse all files
.gitignore
+1
View File
@@ -21,3 +21,4 @@ scripts/COCO_val2017/*
scripts/COCO_val2017.zip
scripts/all_labels.txt
/cmake/cuda_script
/cmake-build-debug/
CMakeLists.txt
+3 -2
View File
@@ -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
README.md
+10 -4
View File
@@ -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)
demo/demo/demo.cpp
+57 -57
View File
@@ -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<std::string>(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<std::string>(conf, "input", "../demo/yolo_test.mp4");
std::string cfgPath = YAMLgetConf<std::string>(conf,"cfg_input", "../tests/darknet/cfg/yolo4tiny.cfg");
std::string namePath = YAMLgetConf<std::string>(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<char>(conf, "ntype", 'y');
int n_classes = YAMLgetConf<int>(conf, "n_classes", 80);
int n_batch = YAMLgetConf<int>(conf, "n_batch", 1);
if(n_batch < 1 || n_batch > 64)
FatalError("Batch dim not supported");
float conf_thresh = YAMLgetConf<float>(conf, "conf_thresh", 0.3);
bool show = YAMLgetConf<bool>(conf, "show", true);
bool save = YAMLgetConf<bool>(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<<COL_GREENB<<"\n\nTime stats:\n";
std::cout<<"Min: "<<*std::min_element(detNN->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; i<detNN->stats.size(); i++) mean += detNN->stats[i]; mean /= detNN->stats.size();
std::cout<<"Avg: "<<mean/n_batch<<" ms\t"<<1000/(mean/n_batch)<<" FPS\n"<<COL_END;
demo/demoConfig.yaml
+22
View File
@@ -0,0 +1,22 @@
# video input
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'
n_classes : 80
n_batch : 1
conf_thresh : 0.3
# demo config
show : true
save : false
docker/Dockerfile.base
+121 -38
View File
@@ -1,57 +1,140 @@
FROM nvidia/cuda:11.3.1-devel-ubuntu20.04
LABEL maintainer "Francesco Gatti"
ENV DEBIAN_FRONTEND=noninteractive
RUN apt-get update && apt-get install 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
RUN DEBIAN_FRONTEND=noninteractive apt-get update && apt install -y git wget libeigen3-dev libyaml-cpp-dev gcc-9 g++-9 libopengl-dev libgl-dev
RUN cd /tmp && \
wget https://github.com/Kitware/CMake/releases/download/v3.21.4/cmake-3.21.4-Linux-x86_64.sh && \
chmod +x cmake-3.21.4-Linux-x86_64.sh && \
./cmake-3.21.4-Linux-x86_64.sh --prefix=/usr/local --exclude-subdir --skip-license && \
rm ./cmake-3.21.4-Linux-x86_64.sh
FROM nvidia/cudagl:11.3.1-devel-ubuntu20.04
LABEL maintainer "TKDNN AUTHORS"
LABEL Description="tkDNN+cudagl"
LABEL com.tkdnn.nvidia.version="11.3.1"
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
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"]
docker/README.md
+1 -4
View File
@@ -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
```
docker/assets/entrypoint_setup.sh
Executable
+123
View File
@@ -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
"$@"
docker/assets/terminator_config
+18
View File
@@ -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"
docker/docker_launch.sh
Executable
+9
View File
@@ -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
docs/demo.md
+19 -24
View File
@@ -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 <path-to-config>
```
In general the demo program takes 1 parameter, the ```<path-to-config>``` 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 <network-rt-file> <path-to-video> <kind-of-network> <number-of-classes> <cfg-path> <name-path> <n-batches> <show-flag> <conf-thresh>
```
where
* ```<network-rt-file>``` is the rt file generated by a test
* ```<<path-to-video>``` is the path to a video file or a camera input
* ```<kind-of-network>``` is the type of network. Thee types are currently supported: ```y``` (YOLO family), ```c``` (CenterNet family) and ```m``` (MobileNet-SSD family)
* ```<number-of-classes>```is the number of classes the network is trained on
* ```<cfg-path> ```is the relative path to the config file (only for darknet based networks) used to train the network
* ```<name-path>```is the relative path to the names file (only for darknet based networks) used to train the network
* ```<n-batches>``` 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).
* ```<show-flag>``` if set to 0 the demo will not show the visualization but save the video into result.mp4 (if n-batches ==1)
* ```<conf-thresh>``` 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.
include/tkDNN/Layer.h
+27 -1
View File
@@ -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:
include/tkDNN/NetworkRT.h
+9 -1
View File
@@ -23,6 +23,8 @@
#include <pluginsRT/ShortcutRT.h>
#include <pluginsRT/UpsampleRT.h>
#include <pluginsRT/YoloRT.h>
#include <pluginsRT/ConstantPaddingRT.h>
#include <pluginsRT/ReflectionPadding.h>
@@ -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();
include/tkDNN/kernels.h
+7
View File
@@ -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
include/tkDNN/pluginsRT/ConstantPaddingRT.h
+109
View File
@@ -0,0 +1,109 @@
//
// Created by perseusdg on 1/7/22.
//
#ifndef _CONSTANTPADDINGRT_PLUGIN_H
#define _CONSTANTPADDINGRT_PLUGIN_H
#include<cassert>
#include <NvInfer.h>
#include <vector>
#include <utils.h>
#include <kernels.h>
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<PluginField> mPluginAttributes;
std::string mPluginNamespace;
};
REGISTER_TENSORRT_PLUGIN(ConstantPaddingRTPluginCreator);
};
#endif //TKDNN_CONSTANTPADDINGRT_H
include/tkDNN/pluginsRT/FlattenConcatRT.h
+5 -1
View File
@@ -1,3 +1,6 @@
#ifndef _FLATTENCONCATRT_PLUGIN_H
#define _FLATTENCONCATRT_PLUGIN_H
#include<cassert>
#include <NvInfer.h>
#include <vector>
@@ -93,4 +96,5 @@ namespace nvinfer1 {
};
REGISTER_TENSORRT_PLUGIN(FlattenConcatRTPluginCreator);
};
};
#endif
include/tkDNN/pluginsRT/ReflectionPadding.h
+101
View File
@@ -0,0 +1,101 @@
#ifndef _REFLECTIONPADDINGRT_PLUGIN_H
#define _REFLECTIONPADDINGRT_PLUGIN_H
#include<cassert>
#include <NvInfer.h>
#include <vector>
#include <utils.h>
#include <kernels.h>
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<PluginField> mPluginAttributes;
std::string mPluginNamespace;
};
REGISTER_TENSORRT_PLUGIN(ReflectionPaddingRTPluginCreator);
};
#endif
include/tkDNN/utils.h
+16 -1
View File
@@ -6,6 +6,8 @@
#include <fstream>
#include <iomanip>
#include <stdlib.h>
#include <yaml-cpp/yaml.h>
#include "cuda.h"
#include "cuda_runtime_api.h"
@@ -16,7 +18,6 @@
#ifdef __linux__
#include <unistd.h>
#endif
#include <ios>
@@ -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: "<<conf_file<<"\n";
return YAML::LoadFile(conf_file);
}
template<typename T>
inline T YAMLgetConf(YAML::Node conf, std::string key, T defaultVal) {
T val = defaultVal;
if(conf && conf[key]) {
val = conf[key].as<T>();
}
return val;
}
#endif //UTILS_H
scripts/checkExecTimes.py
+37
View File
@@ -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')
src/NetworkRT.cpp
+189 -48
View File
@@ -26,15 +26,15 @@ class Logger : public ILogger {
namespace tk { namespace dnn {
std::map<Layer*, nvinfer1::ITensor*>tensors;
std::map<Layer*, nvinfer1::ITensor*>tensors;
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"<<rt_ver<<")\n";
builderRT = createInferBuilder(loggerRT);
std::cout<<"Float16 support: "<<builderRT->platformHasFastFp16()<<"\n";
std::cout<<"Int8 support: "<<builderRT->platformHasFastInt8()<<"\n";
@@ -42,12 +42,12 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
std::cout<<"DLAs: "<<builderRT->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<IInt8EntropyCalibrator> 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; i<net->num_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: "<<builderRT->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<nvinfer1::ICudaEngine>(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"<<std::endl;
builderActive = true;
serialize(name,serializedEngineRT);
delete serializedEngineRT;
#if NV_TENSORRT_MAJOR >= 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 = "<<buf_input_idx<<" -> output index = "<<buf_output_idx<<"\n";
@@ -258,6 +275,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
return convert_layer(input, (Upsample*) l);
if(type == LAYER_DEFORMCONV2D)
return convert_layer(input, (DeformConv2d*) l);
if(type == LAYER_PADDING)
return convert_layer(input, (Padding*) l);
std::cout<<l->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: "<<d.d[0]<<" "<<d.d[1]<<" "<<d.d[2]<<" "<<d.d[3]<<"\n";
}
#else
if(!l->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: "<<d.d[0]<<" "<<d.d[1]<<" "<<d.d[2]<<" "<<d.d[3]<<"\n";
}
#endif
checkNULL(lRT);
if(l->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<<lRT->getNbOutputs()<<std::endl;
IScaleLayer *lRT2 = networkRT->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<ITensor &>(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<PluginField> 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<PluginField> 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<PluginField> 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<<lRT->getNbOutputs()<<std::endl;
IScaleLayer *lRT2 = networkRT->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<const char*>(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;
}
}
src/Padding.cpp
+45
View File
@@ -0,0 +1,45 @@
//
// Created by perseusdg on 03/01/22.
//
#include <iostream>
#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;
}
}}
src/kernels/padding.cu
+110
View File
@@ -0,0 +1,110 @@
#include "kernels.h"
#include <thrust/pair.h>
#include <stdio.h>
/*
* Reflection padding is from https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/native/cuda/ReflectionPad.cu
*/
__device__
inline thrust::pair<int32_t,int32_t> 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<int32_t,int32_t>(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<size_y;block_y += 65535){
int32_t block_y_size = std::min(size_y - block_y,static_cast<int32_t>(65535));
for(int32_t block_z=0;block_z<size_z;block_z += 65535){
int32_t block_z_size = std::min(size_z -block_z,static_cast<int32_t>(65535));
dim3 grid_size(ceilDiv(output_plane_size,static_cast<int32_t>(256)),block_y_size,block_z_size);
reflection_pad2d_out_kernel<<<grid_size,block_size,0,cudaStream>>>(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<int32_t>(256)),c,n);
constant_pad2d_kernel<<<grid_size,block_size,0,cudaStream>>>(srcData,dstData,padT,padL,constant,n,c,input_h,input_w,output_h,output_w);
}
src/pluginsRT/ConstantPaddingRT.cpp
+201
View File
@@ -0,0 +1,201 @@
#include <tkDNN/pluginsRT/ConstantPaddingRT.h>
using namespace nvinfer1;
std::vector<PluginField> 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<const char*>(data),*bufcheck=buf;
padH = readBUF<int32_t>(buf);
padW = readBUF<int32_t>(buf);
i_h = readBUF<int32_t>(buf);
i_w = readBUF<int32_t>(buf);
o_h = readBUF<int32_t>(buf);
o_w = readBUF<int32_t>(buf);
n = readBUF<int32_t>(buf);
c = readBUF<int32_t>(buf);
constant = readBUF<float>(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<const dnnType*>(inputs[0]);
dnnType* dstData = reinterpret_cast<dnnType*>(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<const dnnType*>(inputs[0]);
dnnType* dstData = reinterpret_cast<dnnType*>(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<char*>(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<const int32_t*>(fields[0].data));
int padW = *(static_cast<const int32_t*>(fields[1].data));
int inputH = *(static_cast<const int32_t*>(fields[2].data));
int inputW = *(static_cast<const int32_t*>(fields[3].data));
int outputH = *(static_cast<const int32_t*>(fields[4].data));
int outputW = *(static_cast<const int32_t*>(fields[5].data));
int n = *(static_cast<const int32_t*>(fields[6].data));
int c = *(static_cast<const int32_t*>(fields[7].data));
float constant = *(static_cast<const float*>(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;
}
src/pluginsRT/ReflectionPadding.cpp
+202
View File
@@ -0,0 +1,202 @@
#include <tkDNN/pluginsRT/ReflectionPadding.h>
using namespace nvinfer1;
std::vector<PluginField> 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<const char*>(data),*bufcheck=buf;
padH = readBUF<int32_t>(buf);
padW = readBUF<int32_t>(buf);
input_h = readBUF<int32_t>(buf);
input_w = readBUF<int32_t>(buf);
output_h = readBUF<int32_t>(buf);
output_w = readBUF<int32_t>(buf);
n = readBUF<int32_t>(buf);
c = readBUF<int32_t>(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<const dnnType*>(inputs[0]);
dnnType* dstData = reinterpret_cast<dnnType*>(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<const dnnType*>(inputs[0]);
dnnType* dstData = reinterpret_cast<dnnType*>(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<char*>(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<const int32_t*>(fields[0].data));
int padW = *(static_cast<const int32_t*>(fields[1].data));
int inputH = *(static_cast<const int32_t*>(fields[2].data));
int inputW = *(static_cast<const int32_t*>(fields[3].data));
int outputH = *(static_cast<const int32_t*>(fields[4].data));
int outputW = *(static_cast<const int32_t*>(fields[5].data));
int n = *(static_cast<const int32_t*>(fields[6].data));
int c = *(static_cast<const int32_t*>(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;
}