TensorRT 7.1 Open Source Release

Signed-off-by: Rajeev Rao <rajeevrao@nvidia.com>
This commit is contained in:
Rajeev Rao
2020-07-01 15:14:42 -07:00
parent 2b8863ddff
commit 805810b112
496 changed files with 43444 additions and 6637 deletions
+3
View File
@@ -0,0 +1,3 @@
/.git*
build*
/third_party
+7 -1
View File
@@ -1,2 +1,8 @@
build/
docker/jetpack_files/*
/demo/BERT/models
/demo/BERT/engines
/demo/BERT/squad/*.json
/docker/jetpack_files/*
*.nvmk
*.sln
*.vcxproj
+4 -4
View File
@@ -1,7 +1,3 @@
[submodule "parsers/onnx"]
path = parsers/onnx
url = https://github.com/onnx/onnx-tensorrt.git
branch = 7.0
[submodule "third_party/protobuf"]
path = third_party/protobuf
url = https://github.com/protocolbuffers/protobuf.git
@@ -10,3 +6,7 @@
path = third_party/cub
url = https://github.com/NVlabs/cub.git
branch = 1.8.0
[submodule "parsers/onnx"]
path = parsers/onnx
url = https://github.com/onnx/onnx-tensorrt.git
branch = 7.1
+75 -131
View File
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -19,7 +19,7 @@ include(cmake/modules/set_ifndef.cmake)
include(cmake/modules/find_library_create_target.cmake)
set_ifndef(TRT_LIB_DIR ${CMAKE_BINARY_DIR})
set_ifndef(TRT_BIN_DIR ${CMAKE_BINARY_DIR})
set_ifndef(TRT_OUT_DIR ${CMAKE_BINARY_DIR})
file(STRINGS "${CMAKE_CURRENT_SOURCE_DIR}/include/NvInferVersion.h" VERSION_STRINGS REGEX "#define NV_TENSORRT_.*")
@@ -33,8 +33,8 @@ foreach(TYPE MAJOR MINOR PATCH)
string(REGEX MATCH "[0-9]" TRT_SO_${TYPE} ${TRT_TYPE_STRING})
endforeach(TYPE)
set(TRT_VERSION "${TRT_MAJOR}.${TRT_MINOR}.${TRT_PATCH}.${TRT_BUILD}" CACHE STRING "TRT project version")
set(TRT_SOVERSION "${TRT_SO_MAJOR}.${TRT_SO_MINOR}.${TRT_SO_PATCH}" CACHE STRING "TRT library so version")
set(TRT_VERSION "${TRT_MAJOR}.${TRT_MINOR}.${TRT_PATCH}" CACHE STRING "TensorRT project version")
set(TRT_SOVERSION "${TRT_SO_MAJOR}" CACHE STRING "TensorRT library so version")
message("Building for TensorRT version: ${TRT_VERSION}, library version: ${TRT_SOVERSION}")
if(NOT DEFINED CMAKE_TOOLCHAIN_FILE)
@@ -56,26 +56,81 @@ endif(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT)
option(BUILD_PLUGINS "Build TensorRT plugin" ON)
option(BUILD_PARSERS "Build TensorRT parsers" ON)
option(BUILD_SAMPLES "Build TensorRT samples" ON)
option(NVPARTNER "Build partner repos from source" OFF)
option(NVINTERNAL "Build in NVIDIA internal source tree" OFF)
set(CMAKE_CXX_STANDARD 11)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(CMAKE_CXX_EXTENSIONS OFF)
set(CMAKE_CXX_FLAGS "-Wno-deprecated-declarations ${CMAKE_CXX_FLAGS} -DBUILD_SYSTEM=cmake_oss")
############################# CROSS COMPILATION SETTINGS ##################################
############################################################################################
# Cross-compilation settings
set_ifndef(TRT_PLATFORM_ID "x86_64")
message(STATUS "Targeting TRT Platform: ${TRT_PLATFORM_ID}")
############################################################################################
# Debug settings
set(TRT_DEBUG_POSTFIX _debug CACHE STRING "suffix for debug builds")
if (CMAKE_BUILD_TYPE STREQUAL "Debug")
message("Building in debug mode ${DEBUG_POSTFIX}")
endif()
set(CMAKE_CXX_FLAGS "-Wno-deprecated-declarations ${CMAKE_CXX_FLAGS} -DBUILD_SYSTEM=cmake_oss")
############################################################################################
# Dependencies
set(DEFAULT_CUDA_VERSION 11.0)
set(DEFAULT_CUDNN_VERSION 8.0)
set(DEFAULT_PROTOBUF_VERSION 3.0.0)
set(DEFAULT_CUB_VERSION 1.8.0)
# Dependency Version Resolution
set_ifndef(CUDA_VERSION ${DEFAULT_CUDA_VERSION})
message(STATUS "CUDA version set to ${CUDA_VERSION}")
set_ifndef(CUDNN_VERSION ${DEFAULT_CUDNN_VERSION})
message(STATUS "cuDNN version set to ${CUDNN_VERSION}")
set_ifndef(PROTOBUF_VERSION ${DEFAULT_PROTOBUF_VERSION})
message(STATUS "Protobuf version set to ${PROTOBUF_VERSION}")
find_package(Threads REQUIRED)
if (BUILD_PLUGINS OR BUILD_PARSERS)
include(third_party/zlib.cmake)
include(third_party/protobuf.cmake)
endif()
if(NOT CUB_ROOT_DIR)
set(CUB_ROOT_DIR ${CMAKE_CURRENT_SOURCE_DIR}/third_party/cub CACHE STRING "directory of CUB installation")
endif()
## find_package(CUDA) is broken for cross-compilation. Enable CUDA language instead.
if(NOT DEFINED CMAKE_TOOLCHAIN_FILE)
find_package(CUDA ${CUDA_VERSION} REQUIRED)
endif()
include_directories(
${CUDA_INCLUDE_DIRS}
${CUDNN_ROOT_DIR}/include
)
find_library(CUDNN_LIB cudnn HINTS
${CUDA_TOOLKIT_ROOT_DIR} ${CUDNN_ROOT_DIR} PATH_SUFFIXES lib64 lib)
find_library(CUBLAS_LIB cublas HINTS
${CUDA_TOOLKIT_ROOT_DIR} PATH_SUFFIXES lib64 lib lib/stubs)
find_library(CUBLASLT_LIB cublasLt HINTS
${CUDA_TOOLKIT_ROOT_DIR} PATH_SUFFIXES lib64 lib lib/stubs)
if(BUILD_PARSERS)
configure_protobuf(${PROTOBUF_VERSION})
endif()
find_library_create_target(nvinfer nvinfer SHARED ${TRT_LIB_DIR})
find_library_create_target(nvuffparser nvparsers SHARED ${TRT_LIB_DIR})
find_library(CUDART_LIB cudart HINTS ${CUDA_TOOLKIT_ROOT_DIR} PATH_SUFFIXES lib lib64)
find_library(RT_LIB rt)
set(CUDA_LIBRARIES ${CUDART_LIB})
############################################################################################
# CUDA targets
if (DEFINED GPU_ARCHS)
message(STATUS "GPU_ARCHS defined as ${GPU_ARCHS}. Generating CUDA code for SM ${GPU_ARCHS}")
@@ -88,6 +143,12 @@ else()
70
75
)
if (CUDA_VERSION VERSION_GREATER_EQUAL 11.0)
# Ampere GPU (SM80) support is only available in CUDA versions > 11.0
list(APPEND GPU_ARCHS 80)
else()
message(WARNING "Detected CUDA version is < 11.0. SM80 not supported.")
endif()
message(STATUS "GPU_ARCHS is not defined. Generating CUDA code for default SMs: ${GPU_ARCHS}")
endif()
set(BERT_GENCODES)
@@ -106,140 +167,23 @@ if (${LATEST_SM} GREATER_EQUAL 70)
endif()
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler -Wno-deprecated-declarations")
################################### DEPENDENCIES ##########################################
set(DEFAULT_CUDA_VERSION 10.2)
set(DEFAULT_CUDNN_VERSION 7.6)
set(DEFAULT_PROTOBUF_VERSION 3.0.0)
set(DEFAULT_PROTOBUF_INTERNAL_VERSION 10.0)
set(DEFAULT_CUB_VERSION 1.8.0)
# Dependency Version Resolution
set_ifndef(CUDA_VERSION ${DEFAULT_CUDA_VERSION})
message(STATUS "CUDA version set to ${CUDA_VERSION}")
set_ifndef(CUDNN_VERSION ${DEFAULT_CUDNN_VERSION})
message(STATUS "cuDNN version set to ${CUDNN_VERSION}")
if (NVINTERNAL)
#TODO: Change this to set_ifndef(PROTOBUF_INTERNAL_VERSION ${DEFAULT_PROTOBUF_INTERNAL_VERSION}) once onnx-tensorrts build system is fixed
set_ifndef(PROTOBUF_VERSION ${DEFAULT_PROTOBUF_VERSION})
message(STATUS "Protobuf version set to ${PROTOBUF_INTERNAL_VERSION}")
set_ifndef(CUB_VERSION ${DEFAULT_CUB_VERSION})
message(STATUS "CUB version set to ${CUB_VERSION}")
#TODO: Remove this once CMake is fully intergrated in the P4 build system
set_ifndef(NVINTERNAL_SUFFIX "V2")
else()
set_ifndef(PROTOBUF_VERSION ${DEFAULT_PROTOBUF_VERSION})
message(STATUS "Protobuf version set to ${PROTOBUF_VERSION}")
endif()
find_package(Threads REQUIRED)
if (BUILD_PLUGINS OR BUILD_PARSERS)
include(third_party/zlib.cmake)
include(third_party/protobuf.cmake)
endif()
if (NVINTERNAL)
########################################### DEPENDENCIES FOR BUILDING IN NVIDIA's TREE ############################################
set(EXTERNALS ${PROJECT_SOURCE_DIR}/../externals)
set(CUB_ROOT_DIR ${EXTERNALS}/cub/${CUB_VERSION} CACHE STRING "directory of CUB installation")
set(Protobuf_DIR ${EXTERNALS}/protobuf/${TRT_PLATFORM_ID} CACHE STRING "directory of PROTOBUF installation")
## This needs to be fixed to work with externals
if(NOT DEFINED CMAKE_TOOLCHAIN_FILE)
find_package(CUDA REQUIRED)
endif()
# Set this for ONNX Parser
set(CUDNN_ROOT_DIR ${EXTERNALS}/cudnn/${TRT_PLATFORM_ID}/${CUDNN_VERSION}/cuda-${CUDA_VERSION})
include_directories(
${CUDNN_ROOT_DIR}/include
${CUDA_TOOLKIT_ROOT_DIR}/include
/usr/local/cuda-${CUDA_VERSION}/include
)
#Check externals before using system
find_library(CUDNN_LIB cudnn HINTS
${CUDNN_ROOT_DIR}/lib64 NO_DEFAULT_PATH)
find_library(CUDNN_LIB cudnn HINTS
${CUDNN_ROOT_DIR}/lib64)
find_library(CUBLAS_LIB cublas HINTS
${CUDA_TOOLKIT_ROOT_DIR}/lib NO_DEFAULT_PATH)
find_library(CUBLAS_LIB cublas HINTS
${CUDA_TOOLKIT_ROOT_DIR}/lib)
if(BUILD_PARSERS)
#TODO: Change this to configure_protobuf_internal(${PROTOBUF_INTERNAL_VERSION}) once onnx-tensorrts build system is fixed
configure_protobuf(${PROTOBUF_VERSION})
endif()
########################################### DEPENDENCIES FOR BUILDING IN NVIDIA's TREE ############################################
else()
########################################### DEPENDENCIES FOR BUILDING OUTSIDE OF NVIDIA ############################################
if(NOT CUB_ROOT_DIR)
set(CUB_ROOT_DIR ${CMAKE_CURRENT_SOURCE_DIR}/third_party/cub CACHE STRING "directory of CUB installation")
endif()
## find_package(CUDA) is broken for cross-compilation. Enable CUDA language instead.
if(NOT DEFINED CMAKE_TOOLCHAIN_FILE)
find_package(CUDA ${CUDA_VERSION} REQUIRED)
endif()
include_directories(
${CUDA_INCLUDE_DIRS}
${CUDNN_ROOT_DIR}/include
)
find_library(CUDNN_LIB cudnn HINTS
${CUDA_TOOLKIT_ROOT_DIR} ${CUDNN_ROOT_DIR} PATH_SUFFIXES lib64 lib)
find_library(CUBLAS_LIB cublas HINTS
${CUDA_TOOLKIT_ROOT_DIR} PATH_SUFFIXES lib64 lib lib/stubs)
# CUBLASLT libraries are only available in CUDA versions > 10. Check for CUDA version here and
# remove dependency on the libarary and unset BERT_GENCODES.
if (CUDA_VERSION VERSION_LESS_EQUAL 10.0)
message(WARNING "Detected CUDA version is <= 10.0! Removing BERT plugins from compilation list.")
unset(BERT_GENCODES)
else()
find_library(CUBLASLT_LIB cublasLt HINTS
${CUDA_TOOLKIT_ROOT_DIR} PATH_SUFFIXES lib64 lib lib/stubs)
endif()
if(BUILD_PARSERS)
configure_protobuf(${PROTOBUF_VERSION})
endif()
########################################### DEPENDENCIES FOR BUILDING OUTSIDE OF NVIDIA ############################################
endif()
find_library_create_target(nvinfer nvinfer SHARED ${TRT_LIB_DIR})
if (NOT (NVINTERNAL OR NVPARTNER))
find_library_create_target(nvuffparser nvparsers SHARED ${TRT_LIB_DIR})
endif()
find_library(CUDART_LIB cudart HINTS ${CUDA_TOOLKIT_ROOT_DIR} PATH_SUFFIXES lib lib64)
find_library(RT_LIB rt)
set(CUDA_LIBRARIES ${CUDART_LIB})
############################################################################################
# TensorRT
if(BUILD_PLUGINS)
add_subdirectory(plugin${NVINTERNAL_SUFFIX})
add_subdirectory(plugin)
else()
find_library_create_target(nvinfer_plugin nvinfer_plugin SHARED ${TRT_BIN_DIR} ${TRT_LIB_DIR})
find_library_create_target(nvinfer_plugin nvinfer_plugin SHARED ${TRT_OUT_DIR} ${TRT_LIB_DIR})
endif()
if(BUILD_PARSERS)
add_subdirectory(parsers${NVINTERNAL_SUFFIX})
add_subdirectory(parsers)
else()
if(NVPARTNER OR NVINTERNAL)
find_library_create_target(nvuffparser nvparsers SHARED ${TRT_BIN_DIR} ${TRT_LIB_DIR})
endif()
find_library_create_target(nvcaffeparser nvparsers SHARED ${TRT_BIN_DIR} ${TRT_LIB_DIR})
find_library_create_target(nvonnxparser nvonnxparser SHARED ${TRT_BIN_DIR} ${TRT_LIB_DIR})
find_library_create_target(nvcaffeparser nvparsers SHARED ${TRT_OUT_DIR} ${TRT_LIB_DIR})
find_library_create_target(nvonnxparser nvonnxparser SHARED ${TRT_OUT_DIR} ${TRT_LIB_DIR})
endif()
if(BUILD_SAMPLES)
add_subdirectory(samples${NVINTERNAL_SUFFIX})
add_subdirectory(samples)
endif()
+1 -1
View File
@@ -419,7 +419,7 @@ char const * const errStr = getErrorStr(status);
1. All TensorRT Open Source Software code should contain an NVIDIA copyright header that includes the current year. The following block of text should be prepended to the top of all OSS files. This includes .cpp, .h, .cu, .py, and any other source files which are compiled or interpreted.
```cpp
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
+2 -2
View File
@@ -176,7 +176,7 @@
END OF TERMS AND CONDITIONS
Copyright 2019 NVIDIA Corporation
Copyright 2020 NVIDIA Corporation
Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
@@ -264,4 +264,4 @@
THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
DEALINGS IN THE SOFTWARE.
DEALINGS IN THE SOFTWARE.
+83 -81
View File
@@ -15,8 +15,8 @@ To build the TensorRT OSS components, ensure you meet the following package requ
* [CUDA](https://developer.nvidia.com/cuda-toolkit)
* Recommended versions:
* [cuda-10.2](https://developer.nvidia.com/cuda-10.2-download-archive-base) + cuDNN-7.6
* [cuda-10.0](https://developer.nvidia.com/cuda-10.0-download-archive) + cuDNN-7.6
* cuda-11.0 + cuDNN-8.0
* cuda-10.2 + cuDNN-8.0
* [GNU Make](https://ftp.gnu.org/gnu/make/) >= v4.1
@@ -28,31 +28,39 @@ To build the TensorRT OSS components, ensure you meet the following package requ
* [Python3](https://www.python.org/downloads/release/python-365/) >= v3.6.5
* [PIP](https://pypi.org/project/pip/#history) >= v19.0
* PyPI packages
* [numpy](https://pypi.org/project/numpy/)
* [onnx](https://pypi.org/project/onnx/1.6.0/) 1.6.0
* [onnxruntime](https://pypi.org/project/onnxruntime/) >= 1.3.0
* [pytest](https://pypi.org/project/pytest/)
* Essential libraries and utilities
* [Git](https://git-scm.com/downloads), [pkg-config](https://www.freedesktop.org/wiki/Software/pkg-config/), [Wget](https://www.gnu.org/software/wget/faq.html#download), [Zlib](https://zlib.net/)
* Cross compilation for Jetson platforms requires JetPack's host component installation
* [JetPack](https://developer.nvidia.com/embedded/jetpack) >= 4.2
* [JetPack](https://developer.nvidia.com/embedded/jetpack) >= 4.4
**Optional Packages**
* Containerized builds
* [Docker](https://docs.docker.com/install/) >= 1.12
* [NVIDIA Docker](https://github.com/NVIDIA/nvidia-docker) >= 2.0
* [Docker](https://docs.docker.com/install/) >= 19.03
* [NVIDIA Docker](https://github.com/NVIDIA/nvidia-docker) >= 2.0 or `nvidia-container-toolkit`
* Code formatting tools
* [Clang-format](https://clang.llvm.org/docs/ClangFormat.html)
* [Git-clang-format](https://github.com/llvm-mirror/clang/blob/master/tools/clang-format/git-clang-format)
* Required PyPI packages for Demos
* [Tensorflow-gpu](https://pypi.org/project/tensorflow/1.14.0/) == 1.15.0
**TensorRT Release**
* [TensorRT](https://developer.nvidia.com/nvidia-tensorrt-download) v7.0
* [TensorRT](https://developer.nvidia.com/nvidia-tensorrt-download) v7.1
NOTE: Along with the TensorRT OSS components, the following source packages will also be downloaded, and they are not required to be installed on the system.
- [ONNX-TensorRT](https://github.com/onnx/onnx-tensorrt) v7.0
- [ONNX-TensorRT](https://github.com/onnx/onnx-tensorrt) v7.1
- [CUB](http://nvlabs.github.io/cub/) v1.8.0
- [Protobuf](https://github.com/protocolbuffers/protobuf.git) v3.8.x
@@ -62,7 +70,7 @@ NOTE: Along with the TensorRT OSS components, the following source packages will
1. #### Download TensorRT OSS sources.
```bash
git clone -b master https://github.com/nvidia/TensorRT TensorRT -b release/7.0
git clone -b master https://github.com/nvidia/TensorRT TensorRT
cd TensorRT
git submodule update --init --recursive
export TRT_SOURCE=`pwd`
@@ -70,28 +78,43 @@ NOTE: Along with the TensorRT OSS components, the following source packages will
2. #### Download the TensorRT binary release.
To build the TensorRT OSS, obtain the corresponding TensorRT 7.0 binary release from [NVidia Developer Zone](https://developer.nvidia.com/nvidia-tensorrt-7x-download). For a list of key features, known and fixed issues, refer to the [TensorRT 7.0 Release Notes](https://docs.nvidia.com/deeplearning/sdk/tensorrt-release-notes/tensorrt-7.html#tensorrt-7).
To build the TensorRT OSS, obtain the corresponding TensorRT 7.1 binary release from [NVidia Developer Zone](https://developer.nvidia.com/nvidia-tensorrt-7x-download). For a list of key features, known and fixed issues, refer to the [TensorRT 7.1 Release Notes](https://docs.nvidia.com/deeplearning/tensorrt/release-notes/tensorrt-7.html#rel_7-1-0).
**Example: Ubuntu 18.04 with cuda-10.2**
**Example: Ubuntu 18.04 with cuda-11.0**
Download and extract the latest *TensorRT 7.0 GA package for Ubuntu 18.04 and CUDA 10.2*
Download and extract the latest *TensorRT 7.1 GA package for Ubuntu 18.04 and CUDA 11.0*
```bash
cd ~/Downloads
# Download TensorRT-7.0.0.11.Ubuntu-18.04.x86_64-gnu.cuda-10.2.cudnn7.6.tar.gz
tar -xvzf TensorRT-7.0.0.11.Ubuntu-18.04.x86_64-gnu.cuda-10.2.cudnn7.6.tar.gz
export TRT_RELEASE=`pwd`/TensorRT-7.0.0.11
export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$TRT_RELEASE/lib
tar -xvzf TensorRT-7.1.3.4.Ubuntu-18.04.x86_64-gnu.cuda-11.0.cudnn8.0.tar.gz
export TRT_RELEASE=`pwd`/TensorRT-7.1.3.4
```
**Example: CentOS/RedHat 7 with cuda-10.0**
**Example: Ubuntu 18.04 with cuda-11.0 on PowerPC**
Download and extract the *TensorRT 7.0 GA for CentOS/RedHat 7 and CUDA 10.0 tar package*
Download and extract the latest *TensorRT 7.1 GA package for Ubuntu 18.04 and CUDA 11.0*
```bash
cd ~/Downloads
# Download TensorRT-7.0.0.11.CentOS-7.6.x86_64-gnu.cuda-10.0.cudnn7.6.tar.gz
tar -xvzf TensorRT-7.0.0.11.CentOS-7.6.x86_64-gnu.cuda-10.0.cudnn7.6.tar.gz
export TRT_RELEASE=`pwd`/TensorRT-7.0.0.11
export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$TRT_RELEASE/lib
# Download TensorRT-7.1.3.4.Ubuntu-18.04.powerpc64le-gnu.cuda-11.0.cudnn8.0.tar.gz
tar -xvzf TensorRT-7.1.3.4.Ubuntu-18.04.powerpc64le-gnu.cuda-11.0.cudnn8.0.tar.gz
export TRT_RELEASE=`pwd`/TensorRT-7.1.3.4
```
**Example: CentOS/RedHat 7 with cuda-10.2**
Download and extract the *TensorRT 7.1 GA for CentOS/RedHat 7 and CUDA 10.2 tar package*
```bash
cd ~/Downloads
tar -xvzf TensorRT-7.1.3.4.CentOS-8.0.x86_64-gnu.cuda-10.2.cudnn8.0.tar.gz
export TRT_RELEASE=`pwd`/TensorRT-7.1.3.4
```
**Example: Ubuntu 16.04 with cuda-11.0**
Download and extract the *TensorRT 7.1 GA for Ubuntu 16.04 and CUDA 11.0 tar package*
```bash
cd ~/Downloads
tar -xvzf TensorRT-7.1.3.4.Ubuntu-16.04.x86_64-gnu.cuda-11.0.cudnn8.0.tar.gz
export TRT_RELEASE=`pwd`/TensorRT-7.1.3.4
```
3. #### Download JetPack packages for cross-compilation.[OPTIONAL]
@@ -99,7 +122,7 @@ NOTE: Along with the TensorRT OSS components, the following source packages will
Using the SDK manager, download the host componets of the PDK version or Jetpack specified in the name of the Dockerfile. To do this:
1. [**SDK Manager Step 01**] Log into the SDK manager
2. [**SDK Manager Step 01**] Select the correct platform and Target OS System (should be corresponding to the name of the Dockerfile you are building (e.g. Jetson AGX Xavier, `Linux Jetpack 4.2.1`), then click `Continue`
2. [**SDK Manager Step 01**] Select the correct platform and Target OS System (should be corresponding to the name of the Dockerfile you are building (e.g. Jetson AGX Xavier, `Linux Jetpack 4.4`), then click `Continue`
3. [**SDK Manager Step 02**] Under `Download & Install Options` make note of or change the download folder **and Select Download now. Install later.** then agree to the license terms and click `Continue`
You should now have all expected files to build the container. Move these into the `docker/jetpack_files` folder.
@@ -112,33 +135,44 @@ You should now have all expected files to build the container. Move these into t
1. #### Generate the TensorRT build container.
The docker container can be built using the included Dockerfile. The build container is configured with the environment and packages required for building TensorRT OSS.
The docker container can be built using the included Dockerfiles and build script. The build container is configured with the environment and packages required for building TensorRT OSS.
**Example: Ubuntu 18.04 with cuda-10.2**
**Example: Ubuntu 18.04 with cuda-11.0**
```bash
docker build -f docker/ubuntu.Dockerfile --build-arg UBUNTU_VERSION=18.04 --build-arg CUDA_VERSION=10.2 --tag=tensorrt-ubuntu .
./docker/build.sh --file docker/ubuntu.Dockerfile --tag tensorrt-ubuntu --os 18.04 --cuda 11.0
```
**Example: CentOS/RedHat 7 with cuda-10.0**
**Example: Ubuntu 16.04 with cuda-11.0**
```bash
docker build -f docker/centos.Dockerfile --build-arg CENTOS_VERSION=7 --build-arg CUDA_VERSION=10.0 --tag=tensorrt-centos .
./docker/build.sh --file docker/ubuntu.Dockerfile --tag tensorrt-ubuntu1604 --os 16.04 --cuda 11.0
```
**Example: Cross compile for JetPack 4.2.1 with cuda-10.0**
**Example: CentOS/RedHat 7 with cuda-10.2**
```bash
./docker/build.sh --file docker/centos.Dockerfile --tag tensorrt-centos --os 7 --cuda 10.2
```
**Example: Cross compile for JetPack 4.4 with cuda-10.2**
```bash
docker build -f docker/ubuntu-cross-aarch64.Dockerfile --build-arg UBUNTU_VERSION=18.04 --build-arg CUDA_VERSION=10.0 --tag tensorrt-ubuntu-aarch64 .
`
./docker/build.sh --file docker/ubuntu-cross-aarch64.Dockerfile --tag tensorrt-ubuntu-jetpack --os 18.04 --cuda 10.2
```
**Example: Cross compile for PowerPC with cuda-11.0**
```bash
./docker/build.sh --file docker/ubuntu-cross-ppc64le.Dockerfile --tag tensorrt-ubuntu-ppc --os 18.04 --cuda 11.0
```
2. #### Launch the TensorRT build container.
```bash
docker run -v $TRT_RELEASE:/tensorrt -v $TRT_SOURCE:/workspace/TensorRT -it tensorrt-ubuntu:latest
./docker/launch.sh --tag tensorrt-ubuntu --gpus all --release $TRT_RELEASE --source $TRT_SOURCE
```
> NOTE: To run TensorRT/CUDA programs within the build container, install [nvidia-docker](#prerequisites). Replace the `docker run` command with `nvidia-docker run` or `docker run --runtime=nvidia`.
> NOTE: To run TensorRT/CUDA programs in the build container, install [NVIDIA Docker support](#prerequisites). Docker versions < 19.03 require `nvidia-docker2` and `--runtime=nvidia` flag for docker run commands. On versions >= 19.03, you need the `nvidia-container-toolkit` package and `--gpus <NUM_GPUS>` flag.
## Building The TensorRT OSS Components
@@ -148,28 +182,30 @@ You should now have all expected files to build the container. Move these into t
```bash
cd $TRT_SOURCE
mkdir -p build && cd build
cmake .. -DTRT_LIB_DIR=$TRT_RELEASE/lib -DTRT_BIN_DIR=`pwd`/out
cmake .. -DTRT_LIB_DIR=$TRT_RELEASE/lib -DTRT_OUT_DIR=`pwd`/out
make -j$(nproc)
```
> NOTE:
> 1. The default CUDA version used by CMake is 10.2. To override this, for example to 10.0, append `-DCUDA_VERSION=10.0` to the cmake command.
> 1. The default CUDA version used by CMake is 11.0. To override this, for example to 10.2, append `-DCUDA_VERSION=10.2` to the cmake command.
> 2. Samples may fail to link on CentOS7. To work around this create the following symbolic link:
> `ln -s $TRT_BIN_DIR/libnvinfer_plugin.so $TRT_BIN_DIR/libnvinfer_plugin.so.7`
> `ln -s $TRT_OUT_DIR/libnvinfer_plugin.so $TRT_OUT_DIR/libnvinfer_plugin.so.7`
The required CMake arguments are:
- `TRT_LIB_DIR`: Path to the TensorRT installation directory containing libraries.
- `TRT_BIN_DIR`: Output directory where generated build artifacts will be copied.
- `TRT_OUT_DIR`: Output directory where generated build artifacts will be copied.
The following CMake build parameters are *optional*:
- `CMAKE_BUILD_TYPE`: Specify if binaries generated are for release or debug (contain debug symbols). Values consists of [`Release`] | `Debug`
- `CUDA_VERISON`: The version of CUDA to target, for example [`10.2`].
- `CUDA_VERISON`: The version of CUDA to target, for example [`11.0`].
- `CUDNN_VERSION`: The version of cuDNN to target, for example [`7.6`].
- `CUDNN_VERSION`: The version of cuDNN to target, for example [`8.0`].
- `NVCR_SUFFIX`: Optional nvcr/cuda image suffix. Set to "-rc" for CUDA11 RC builds until general availability. Blank by default.
- `PROTOBUF_VERSION`: The version of Protobuf to use, for example [`3.8.x`]. Note: Changing this will not configure CMake to use a system version of Protobuf, it will configure CMake to download and try building that version.
@@ -183,61 +219,27 @@ You should now have all expected files to build the container. Move these into t
Other build options with limited applicability:
- `NVINTERNAL`: Used by TensorRT team for internal builds. Values consists of [`OFF`] | `ON`.
- `PROTOBUF_INTERNAL_VERSION`: The version of protobuf to use, for example [`10.0`]. Only applicable if `NVINTERNAL` is also enabled.
- `NVPARTNER`: For use by NVIDIA partners with exclusive source access. Values consists of [`OFF`] | `ON`.
- `CUB_VERSION`: The version of CUB to use, for example [`1.8.0`].
- `GPU_ARCHS`: GPU (SM) architectures to target. By default we generate CUDA code for all major SMs. Specific SM versions can be specified here as a quoted space-separated list to reduce compilation time and binary size. Table of compute capabilities of NVIDIA GPUs can be found [here](https://developer.nvidia.com/cuda-gpus). Examples:
- Titan V: `-DGPU_ARCHS="70"`
- Tesla V100: `-DGPU_ARCHS="70"`
- GeForce RTX 2080: `-DGPU_ARCHS="75"`
- Tesla T4: `-DGPU_ARCHS="75"`
- Multiple SMs: `-DGPU_ARCHS="70 75"`
- NVidia A100: `-DGPU_ARCHS="80"`
- Tesla T4, GeForce RTX 2080: `-DGPU_ARCHS="75"`
- Titan V, Tesla V100: `-DGPU_ARCHS="70"`
- Multiple SMs: `-DGPU_ARCHS="80 75"`
## Install the TensorRT OSS Components [Optional]
* Copy the build artifacts into the TensorRT installation directory, updating the installation.
* TensorRT installation directory is determined as `$TRT_LIB_DIR/..`
* Installation might require superuser privileges depending on the path and permissions of files being replaced.
* Installation is not supported in cross compilation scenario. Please copy the result files from `build/out` folder into the target device.
```bash
sudo make install
```
* Verify the TensorRT samples have been installed correctly.
```bash
cd $TRT_LIB_DIR/../bin/
./sample_googlenet
```
If the sample was installed correctly, the following information will be printed out in the terminal.
```bash
[08/23/2019-22:08:57] [I] Building and running a GPU inference engine for GoogleNet
[08/23/2019-22:08:59] [I] [TRT] Some tactics do not have sufficient workspace memory to run. Increasing workspace size may increase performance, please check verbose output.
[08/23/2019-22:09:05] [I] [TRT] Detected 1 inputs and 1 output network tensors.
[08/23/2019-22:09:05] [I] Ran /tensorrt/bin/sample_googlenet with:
[08/23/2019-22:09:05] [I] Input(s): data
[08/23/2019-22:09:05] [I] Output(s): prob
&&&& PASSED TensorRT.sample_googlenet # /tensorrt/bin/sample_googlenet
```
## Useful Resources
#### TensorRT
* [TensorRT Homepage](https://developer.nvidia.com/tensorrt)
* [TensorRT Developer Guide](https://docs.nvidia.com/deeplearning/sdk/tensorrt-developer-guide/index.html)
* [TensorRT Sample Support Guide](https://docs.nvidia.com/deeplearning/sdk/tensorrt-sample-support-guide/index.html)
* [TensorRT Developer Guide](https://docs.nvidia.com/deeplearning/tensorrt/developer-guide/index.html)
* [TensorRT Sample Support Guide](https://docs.nvidia.com/deeplearning/tensorrt/sample-support-guide/index.html)
* [TensorRT Discussion Forums](https://devtalk.nvidia.com/default/board/304/tensorrt/)
## Known Issues
#### TensorRT 7.0
* See [Release Notes](https://docs.nvidia.com/deeplearning/sdk/tensorrt-release-notes/tensorrt-7.html#tensorrt-7).
#### TensorRT 7.1
* [demo/BERT](demo/BERT) has a known accuracy regression for Volta GPUs; F1 score dropped (from 90 in TensorRT 7.0) to 85. A fix is underway.
* See [Release Notes](https://docs.nvidia.com/deeplearning/tensorrt/release-notes/tensorrt-7.html#rel_7-1-3).
+1 -1
View File
@@ -1 +1 @@
7.0.0.11
7.1.3.4
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
+1 -1
View File
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
+3 -13
View File
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -31,12 +31,7 @@ set(CMAKE_CXX_COMPILER_TARGET aarch64)
set(CMAKE_C_COMPILER_FORCED TRUE)
set(CMAKE_CXX_COMPILER_FORCED TRUE)
if(NVINTERNAL)
set(EXT_PATH ${PROJECT_SOURCE_DIR}/../externals)
set(CUDA_ROOT ${EXT_PATH}/cuda-${CUDA_VERSION}-${TRT_PLATFORM_ID}/${CUDA_PLATFORM_ID})
else()
set(CUDA_ROOT /usr/local/cuda-${CUDA_VERSION}/targets/${CUDA_PLATFORM_ID} CACHE STRING "CUDA ROOT dir")
endif()
set(CUDA_ROOT /usr/local/cuda-${CUDA_VERSION}/targets/${CUDA_PLATFORM_ID} CACHE STRING "CUDA ROOT dir")
set(CUDA_TOOLKIT_ROOT_DIR ${CUDA_ROOT})
set(CUDA_INCLUDE_DIRS ${CUDA_ROOT}/include)
@@ -47,12 +42,7 @@ set(CMAKE_CUDA_HOST_COMPILER ${CMAKE_CXX_COMPILER} CACHE STRING "" FORCE)
set(CMAKE_CUDA_FLAGS "-cudart none -I${CUDA_INCLUDE_DIRS} -Xcompiler=\"-fPIC ${CMAKE_CXX_FLAGS}\"" CACHE STRING "" FORCE)
set(CMAKE_CUDA_COMPILER_FORCED TRUE)
if(DEFINED ENV{VULCAN} AND NOT $ENV{VULCAN} STREQUAL "")
message("cmake_aarch64.toolchain using VULCAN mode")
set(CUDA_LIBS ${DEVLIBPATHS} -L${CUDA_ROOT}/lib64)
else()
set(CUDA_LIBS -L${CUDA_ROOT}/lib)
endif()
set(CUDA_LIBS -L${CUDA_ROOT}/lib)
set(ADDITIONAL_PLATFORM_LIB_FLAGS ${CUDA_LIBS} -lcublas -lcudart -lstdc++ -lm)
+1 -4
View File
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -33,7 +33,4 @@ endif()
set(CUDA_INCLUDE_DIRS ${CUDA_ROOT}/include)
if(DEFINED ENV{VULCAN} AND NOT $ENV{VULCAN} STREQUAL "")
set(DISABLE_SWIG TRUE)
endif()
set(TRT_PLATFORM_ID "ppc64le")
+2 -2
View File
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -21,7 +21,7 @@ if(DEFINED ENV{QNX_BASE})
set(QNX_BASE $ENV{QNX_BASE})
message(STATUS "Found QNX_BASE = ${QNX_BASE}")
elseif(DEFINED ENV{TOOLS_BASE})
set(QNX_BASE $ENV{TOOLS_BASE}/embedded/qnx/qnx700-ga3)
set(QNX_BASE $ENV{TOOLS_BASE}/embedded/qnx/qnx700-ga4)
message(STATUS "Found QNX_BASE = ${QNX_BASE}")
else()
message(FATAL_ERROR "QNX_BASE was not found")
+3 -3
View File
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -40,9 +40,9 @@ set(W10_LINKER ${MSVC_COMPILER_DIR}/bin/amd64/link)
set(CMAKE_CUDA_HOST_COMPILER ${CMAKE_NVCC_COMPILER} CACHE STRING "" FORCE)
set(ADDITIONAL_PLATFORM_INCL_FLAGS "-I${MSVC_COMPILER_DIR}/include -I${MSVC_COMPILER_DIR}/../ucrt/include")
set(ADDITIONAL_PLATFORM_LIB_FLAGS ${ADDITIONAL_PLATFORM_LIB_FLAGS} "-LIBPATH:${NV_TOOLS}/ddk/wddmv2/dev/rs4/17130/Lib/10.0.17130.0/um/x64")
set(ADDITIONAL_PLATFORM_LIB_FLAGS ${ADDITIONAL_PLATFORM_LIB_FLAGS} "-LIBPATH:${NV_TOOLS}/ddk/wddmv2/official/17134/Lib/10.0.17134.0/um/x64")
set(ADDITIONAL_PLATFORM_LIB_FLAGS ${ADDITIONAL_PLATFORM_LIB_FLAGS} "-LIBPATH:${MSVC_COMPILER_DIR}/lib/amd64" )
set(ADDITIONAL_PLATFORM_LIB_FLAGS ${ADDITIONAL_PLATFORM_LIB_FLAGS} "-LIBPATH:${MSVC_COMPILER_DIR}/../ucrt/lib/x64")
set(ADDITIONAL_PLATFORM_LIB_FLAGS ${ADDITIONAL_PLATFORM_LIB_FLAGS} "-LIBPATH:${W10_CUDA_ROOT}/lib/x64 cudart.lib cublas.lib")
set(TRT_PLATFORM_ID "win10")
set(TRT_PLATFORM_ID "win10")
+1 -5
View File
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -26,8 +26,4 @@ endif()
set(CUDA_INCLUDE_DIRS ${CUDA_ROOT}/include)
if(DEFINED ENV{VULCAN} AND NOT $ENV{VULCAN} STREQUAL "")
set(DISABLE_SWIG TRUE)
endif()
set(TRT_PLATFORM_ID "x86_64")
+72
View File
@@ -0,0 +1,72 @@
#
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
#
cmake_minimum_required(VERSION 3.12 FATAL_ERROR)
project(infer_c LANGUAGES CXX)
find_package(CUDA)
include(FetchContent)
FetchContent_Declare(
pybind11
GIT_REPOSITORY https://github.com/pybind/pybind11
GIT_TAG v2.2.3
)
FetchContent_GetProperties(pybind11)
if(NOT pybind11_POPULATED)
FetchContent_Populate(pybind11)
add_subdirectory(${pybind11_SOURCE_DIR} ${pybind11_BINARY_DIR})
endif()
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-deprecated-declarations")
include($ENV{TRT_SOURCE}/cmake/modules/set_ifndef.cmake)
set_ifndef(TRT_INC_DIR $ENV{TRT_SOURCE}/include)
set_ifndef(TRT_LIB_DIR $ENV{TRT_RELEASE}/lib)
set_ifndef(TRT_OUT_DIR $ENV{TRT_SOURCE}/build/out)
include_directories(
infer_c
${CUDA_INCLUDE_DIRS}
${TRT_INC_DIR}
)
link_directories(
${TRT_OUT_DIR}
${TRT_LIB_DIR}
)
pybind11_add_module(infer_c
infer_c/infer_c.cpp
infer_c/logging.cpp
)
target_link_libraries(infer_c PRIVATE
${CUDA_LIBRARIES}
nvinfer
nvinfer_plugin
)
add_executable(perf
infer_c/perf.cpp
infer_c/logging.cpp
)
target_link_libraries(perf
${CUDA_LIBRARIES}
nvinfer
nvinfer_plugin
)
+501
View File
@@ -0,0 +1,501 @@
# BERT Inference Using TensorRT
This subfolder of the BERT TensorFlow repository, tested and maintained by NVIDIA, provides scripts to perform high-performance inference using NVIDIA TensorRT.
## Table Of Contents
- [Model Overview](#model-overview)
* [Model Architecture](#model-architecture)
* [TensorRT Inference Pipeline](#tensorrt-inference-pipeline)
* [Version Info](#version-info)
- [Setup](#setup)
* [Requirements](#requirements)
- [Quick Start Guide](#quick-start-guide)
* [(Optional) Trying a different configuration](#optional-trying-a-different-configuration)
- [Advanced](#advanced)
* [Scripts and sample code](#scripts-and-sample-code)
* [Command-line options](#command-line-options)
* [TensorRT inference process](#tensorrt-inference-process)
- [Accuracy](#accuracy)
* [Evaluating Post-Training-Quantization INT8 accuracy](#evaluating-ptq-post-training-quantization-int8-accuracy-using-the-squad-dataset)
* [Evaluating Quantization-Aware-Training INT8 accuracy](#evaluating-qat-quantization-aware-training-int8-accuracy-using-the-squad-dataset)
- [Performance](#performance)
* [Benchmarking](#benchmarking)
* [TensorRT inference benchmark](#tensorrt-inference-benchmark)
* [Results](#results)
* [Inference performance: NVIDIA A100](#inference-performance-nvidia-a100-40gb)
* [BERT Base](#bert-base)
* [BERT Large](#bert-large)
* [Inference performance: NVIDIA T4](#inference-performance-nvidia-t4-16gb)
* [BERT Base](#bert-base-1)
* [BERT Large](#bert-large-1)
* [Inference performance: NVIDIA V100](#inference-performance-nvidia-v100-16gb)
* [BERT Base](#bert-base-2)
* [BERT Large](#bert-large-2)
## Model overview
BERT, or Bidirectional Encoder Representations from Transformers, is a new method of pre-training language representations which obtains state-of-the-art results on a wide array of Natural Language Processing (NLP) tasks. This model is based on the [BERT: Pre-training of Deep Bidirectional Transformers for Language Understanding](https://arxiv.org/abs/1810.04805) paper. NVIDIA's BERT is an optimized version of [Google's official implementation](https://github.com/google-research/bert), leveraging mixed precision arithmetic and Tensor Cores for faster inference times while maintaining target accuracy.
Other publicly available implementations of BERT include:
1. [NVIDIA PyTorch](https://github.com/NVIDIA/DeepLearningExamples/tree/master/PyTorch/LanguageModeling/BERT)
2. [Hugging Face](https://github.com/huggingface/pytorch-pretrained-BERT)
3. [codertimo](https://github.com/codertimo/BERT-pytorch)
4. [gluon-nlp](https://github.com/dmlc/gluon-nlp/tree/master/scripts/bert)
5. [Google's official implementation](https://github.com/google-research/bert)
### Model architecture
BERT's model architecture is a multi-layer bidirectional Transformer encoder. Based on the model size, we have the following two default configurations of BERT:
| **Model** | **Hidden layers** | **Hidden unit size** | **Attention heads** | **Feed-forward filter size** | **Max sequence length** | **Parameters** |
|:---------:|:----------:|:----:|:---:|:--------:|:---:|:----:|
|BERT-Base |12 encoder| 768| 12|4 x 768|512|110M|
|BERT-Large|24 encoder|1024| 16|4 x 1024|512|330M|
Typically, the language model is followed by a few task-specific layers. The model used here includes layers for question answering.
### TensorRT Inference Pipeline
BERT inference consists of three main stages: tokenization, the BERT model, and finally a projection of the tokenized prediction onto the original text.
Since the tokenizer and projection of the final predictions are not nearly as compute-heavy as the model itself, we run them on the host. The BERT model is GPU-accelerated via TensorRT.
The tokenizer splits the input text into tokens that can be consumed by the model. For details on this process, see [this tutorial](https://mccormickml.com/2019/05/14/BERT-word-embeddings-tutorial/).
To run the BERT model in TensorRT, we construct the model using TensorRT APIs and import the weights from a pre-trained TensorFlow checkpoint from [NGC](https://ngc.nvidia.com/models/nvidian:bert_tf_v2_large_fp16_128). Finally, a TensorRT engine is generated and serialized to the disk. The various inference scripts then load this engine for inference.
Lastly, the tokens predicted by the model are projected back to the original text to get a final result.
### Version Info
The following software version configuration has been tested:
|Software|Version|
|--------|-------|
|Python|3.6.9|
|TensorRT|7.1.3.4|
|CUDA|11.0.171|
## Setup
The following section lists the requirements that you need to meet in order to run the BERT model.
### Requirements
This demo BERT application can be run within the TensorRT Open Source build container. If running in a different environment, ensure you have the following packages installed.
* [NGC CLI](https://ngc.nvidia.com/setup/installers/cli) - for downloading BERT checkpoints from NGC.
* PyPI Packages:
* [pycuda](https://pypi.org/project/pycuda/) 2019.1.2
* [onnx](https://pypi.org/project/onnx/1.6.0/) 1.6.0
* [tensorflow](https://pypi.org/project/tensorflow/1.15.3/) 1.15
* NVIDIA [Volta](https://www.nvidia.com/en-us/data-center/volta-gpu-architecture/), [Turing](https://www.nvidia.com/en-us/geforce/turing/) or [Ampere](https://www.nvidia.com/en-us/data-center/nvidia-ampere-gpu-architecture/) based GPU with NVIDIA Driver 450.37 or later.
## Quick Start Guide
1. Build and launch the TensorRT-OSS build container. On x86 with Ubuntu 18.04 for example:
```bash
cd <TensorRT-OSS>
./docker/build.sh --file docker/ubuntu.Dockerfile --tag tensorrt-ubuntu --os 18.04 --cuda 11.0
./docker/launch.sh --tag tensorrt-ubuntu --gpus all --release $TRT_RELEASE --source $TRT_SOURCE
```
**Note:** After this point, all commands should be run from within the container.
2. Build the TensorRT Plugins library from source and install the TensorRT python bindings:
```bash
cd $TRT_SOURCE
export LD_LIBRARY_PATH=`pwd`/build/out:$LD_LIBRARY_PATH:/tensorrt/lib
mkdir -p build && cd build
cmake .. -DTRT_LIB_DIR=$TRT_RELEASE/lib -DTRT_OUT_DIR=`pwd`/out
make -j$(nproc)
pip3 install /tensorrt/python/tensorrt-7.1*-cp36-none-linux_x86_64.whl
```
**Note:** While the workflow and Performance Data presented here are based on plugin library built from source, the BERT sample is also expected to work with pre-compiled libraries shipped with TensorRT releases.
3. Download the SQuAD dataset and BERT checkpoints:
```bash
cd $TRT_SOURCE/demo/BERT
```
Download SQuAD v1.1 training and dev dataset.
```bash
bash scripts/download_squad.sh
```
Download Tensorflow checkpoints for BERT large model with sequence length 128 and fp16 weights, fine-tuned for SQuAD v2.0.
```bash
bash scripts/download_model.sh
````
**Note:** Since the datasets and checkpoints are stored in the directory mounted from the host, they do *not* need to be downloaded each time the container is launched.
4. Build a TensorRT engine. To build an engine, run the `builder.py` script. For example:
```bash
mkdir -p /workspace/TensorRT/demo/BERT/engines && python3 builder.py -m /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_128_v2/model.ckpt-8144 -o /workspace/TensorRT/demo/BERT/engines/bert_large_128.engine -b 1 -s 128 --fp16 -c /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_128_v2
```
This will build an engine with a maximum batch size of 1 (`-b 1`), and sequence length of 128 (`-s 128`) using mixed precision (`--fp16`) using the BERT Large V2 FP16 Sequence Length 128 checkpoint (`-c /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_128_v2`).
5. Run inference. Two options are provided for running the model.
a. `inference.py` script
This script accepts a passage and question and then runs the engine to generate an answer.
For example:
```bash
python3 inference.py -e /workspace/TensorRT/demo/BERT/engines/bert_large_128.engine -p "TensorRT is a high performance deep learning inference platform that delivers low latency and high throughput for apps such as recommenders, speech and image/video on NVIDIA GPUs. It includes parsers to import models, and plugins to support novel ops and layers before applying optimizations for inference. Today NVIDIA is open-sourcing parsers and plugins in TensorRT so that the deep learning community can customize and extend these components to take advantage of powerful TensorRT optimizations for your apps." -q "What is TensorRT?" -v /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_128_v2/vocab.txt
```
b. `inference.ipynb` Jupyter Notebook
The Jupyter Notebook includes a passage and various example questions and allows you to interactively make modifications and see the outcome.
To launch the Jupyter Notebook from inside the container, run:
```bash
jupyter notebook --ip 0.0.0.0 inference.ipynb
```
Then, use your browser to open the link displayed. The link should look similar to: `http://127.0.0.1:8888/?token=<TOKEN>`
6. Run inference with CUDA Graph support.
A separate python `inference_c.py` script is provided to run inference with CUDA Graph support. This is necessary since CUDA Graph is only supported through CUDA C/C++ APIs, not pyCUDA. The `inference_c.py` script uses pybind11 to interface with C/C++ for CUDA graph capturing and launching. The cmdline interface is the same as `inference.py` except for an extra `--enable-graph` option.
```bash
mkdir -p build
cd build; cmake ..
make; cd ..
python3 inference_c.py -e /workspace/TensorRT/demo/BERT/engines/bert_large_128.engine --enable-graph -p "TensorRT is a high performance deep learning inference platform that delivers low latency and high throughput for apps such as recommenders, speech and image/video on NVIDIA GPUs. It includes parsers to import models, and plugins to support novel ops and layers before applying optimizations for inference. Today NVIDIA is open-sourcing parsers and plugins in TensorRT so that the deep learning community can customize and extend these components to take advantage of powerful TensorRT optimizations for your apps." -q "What is TensorRT?" -v /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_128_v2/vocab.txt
```
A separate C/C++ inference benchmark executable `perf` (compiled from `perf.cpp`) is provided to run inference benchmarks with CUDA Graph. The cmdline interface is the same as `perf.py` except for an extra `--enable_graph` option.
```bash
build/perf -e /workspace/TensorRT/demo/BERT/engines/bert_large_128.engine -b 1 -s 128 -w 100 -i 1000 --enable_graph
```
### (Optional) Trying a different configuration
If you would like to run another configuration, you can manually download checkpoints using the included script. For example, run:
```bash
bash scripts/download_model.sh base
```
to download a BERT Base model instead of the default BERT Large model.
To view all available model options, run:
```bash
bash scripts/download_model.sh -h
```
## Advanced
The following sections provide greater details on inference with TensorRT.
### Scripts and sample code
In the `root` directory, the most important files are:
- `builder.py` - Builds an engine for the specified BERT model
- `Dockerfile` - Container which includes dependencies and model checkpoints to run BERT
- `inference.ipynb` - Runs inference interactively
- `inference.py` - Runs inference with a given passage and question
- `perf.py` - Runs inference benchmarks
The `scripts/` folder encapsulates all the one-click scripts required for running various supported functionalities, such as:
- `build.sh` - Builds a Docker container that is ready to run BERT
- `launch.sh` - Launches the container created by the `build.sh` script.
- `download_model.sh` - Downloads pre-trained model checkpoints from NGC
- `inference_benchmark.sh` - Runs an inference benchmark and prints results
Other folders included in the `root` directory are:
- `helpers` - Contains helpers for tokenization of inputs
The `infer_c/` folder contains all the necessary C/C++ files required for CUDA Graph support.
- `bert_infer.h` - Defines necessary data structures for running BERT inference
- `infer_c.cpp` - Defines C/C++ interface using pybind11 that can be plugged into `inference_c.py`
- `perf.cpp` - Runs inference benchmarks. It is equivalent to `perf.py`, with an extra option `--enable_graph` to enable CUDA Graph support.
### Command-line options
To view the available parameters for each script, you can use the help flag (`-h`).
### TensorRT inference process
As mentioned in the [Quick Start Guide](#quick-start-guide), two options are provided for running inference:
1. The `inference.py` script which accepts a passage and a question and then runs the engine to generate an answer. Alternatively, this script can be used to run inference on the Squad dataset.
2. The `inference.ipynb` Jupyter Notebook which includes a passage and various example questions and allows you to interactively make modifications and see the outcome.
## Accuracy
### Evaluating PTQ (post-training quantization) Int8 Accuracy Using The SQuAD Dataset
1. Download Tensorflow checkpoints for a BERT Large FP16 SQuAD v2 model with a sequence length of 384:
```bash
bash scripts/download_model.sh large fp16 384 v2
```
2. Build an engine:
**Turing and Ampere GPUs**
```bash
# QKVToContextPlugin and SkipLayerNormPlugin supported with INT8 I/O. To enable, use -imh and -iln builder flags respectively.
mkdir -p /workspace/TensorRT/demo/BERT/engines && python3 builder.py -m /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2/model.ckpt-8144 -o /workspace/TensorRT/demo/BERT/engines/bert_large_384_int8mix.engine -b 1 -s 384 --int8 --fp16 --strict -c /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2 --squad-json ./squad/train-v1.1.json -v /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2/vocab.txt --calib-num 100 -iln -imh
```
**Xavier GPU**
```bash
# Only supports SkipLayerNormPlugin running with INT8 I/O. Use -iln builder flag to enable.
mkdir -p /workspace/TensorRT/demo/BERT/engines && python3 builder.py -m /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2/model.ckpt-8144 -o /workspace/TensorRT/demo/BERT/engines/bert_large_384_int8mix.engine -b 1 -s 384 --int8 --fp16 --strict -c /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2 --squad-json ./squad/train-v1.1.json -v /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2/vocab.txt --calib-num 100 -iln
```
**Volta GPU**
```bash
# No support for QKVToContextPlugin or SkipLayerNormPlugin running with INT8 I/O. Don't specify -imh or -iln in builder flags.
mkdir -p /workspace/TensorRT/demo/BERT/engines && python3 builder.py -m /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2/model.ckpt-8144 -o /workspace/TensorRT/demo/BERT/engines/bert_large_384_int8mix.engine -b 1 -s 384 --int8 --fp16 --strict -c /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2 --squad-json ./squad/train-v1.1.json -v /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2/vocab.txt --calib-num 100
```
This will build an engine with a maximum batch size of 1 (`-b 1`), calibration dataset squad (`--squad-json ./squad/train-v1.1.json`), calibration sentences number 100 (`--calib-num 100`), and sequence length of 384 (`-s 384`) using INT8 mixed precision computation where possible (`--int8 --fp16 --strict`).
3. Run inference using the squad dataset, and evaluate the F1 score and exact match score:
```bash
python3 inference.py -e /workspace/TensorRT/demo/BERT/engines/bert_large_384_int8mix.engine -s 384 -sq ./squad/dev-v1.1.json -v /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2/vocab.txt -o ./predictions.json
python3 squad/evaluate-v1.1.py squad/dev-v1.1.json ./predictions.json 90
```
### Evaluating QAT (quantization aware training) Int8 Accuracy Using The SQuAD Dataset
1. Download checkpoint for BERT Large FP16 SQuAD v1.1 model with sequence length of 384:
```bash
bash scripts/download_model.sh pyt v1_1
```
2. Build an engine:
**Turing and Ampere GPUs**
```bash
# QKVToContextPlugin and SkipLayerNormPlugin supported with INT8 I/O. To enable, use -imh and -iln builder flags respectively.
mkdir -p /workspace/TensorRT/demo/BERT/engines && python3 builder.py -o /workspace/TensorRT/demo/BERT/engines/bert_large_384_int8mix.engine -b 1 -s 384 --int8 --fp16 --strict -c /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2 -v /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2/vocab.txt -x /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_pyt_onnx_large_qa_squad11_amp_fake_quant_v1/bert_large_v1_1_fake_quant.onnx -iln -imh
```
**Xavier GPU**
```bash
# Only supports SkipLayerNormPlugin running with INT8 I/O. Use -iln builder flag to enable.
mkdir -p /workspace/TensorRT/demo/BERT/engines && python3 builder.py -o /workspace/TensorRT/demo/BERT/engines/bert_large_384_int8mix.engine -b 1 -s 384 --int8 --fp16 --strict -c /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2 -v /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2/vocab.txt -x /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_pyt_onnx_large_qa_squad11_amp_fake_quant_v1/bert_large_v1_1_fake_quant.onnx -iln
```
**Volta GPU**
```bash
# No support for QKVToContextPlugin or SkipLayerNormPlugin running with INT8 I/O. Don't specify -imh or -iln in builder flags.
mkdir -p /workspace/TensorRT/demo/BERT/engines && python3 builder.py -o /workspace/TensorRT/demo/BERT/engines/bert_large_384_int8mix.engine -b 1 -s 384 --int8 --fp16 --strict -c /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2 -v /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2/vocab.txt -x /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_pyt_onnx_large_qa_squad11_amp_fake_quant_v1/bert_large_v1_1_fake_quant.onnx
```
This will build and engine with a maximum batch size of 1 (`-b 1`) and sequence length of 384 (`-s 384`) using INT8 mixed precision computation where possible (`--int8 --fp16 --strict`).
3. Run inference using the squad dataset, and evaluate the F1 score and exact match score:
```bash
python3 inference.py -e /workspace/TensorRT/demo/BERT/engines/bert_large_384_int8mix.engine -s 384 -sq ./squad/dev-v1.1.json -v /workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_large_fp16_384_v2/vocab.txt -o ./predictions.json
python3 squad/evaluate-v1.1.py squad/dev-v1.1.json ./predictions.json 90
```
## Performance
### Benchmarking
The following section shows how to run benchmarks measuring the model performance in inference modes.
#### TensorRT inference benchmark
The inference benchmark is performed on a single GPU by the `inference_benchmark.sh` script, which takes the following steps for each set of model parameters:
1. Downloads checkpoints and builds a TensorRT engine if it does not already exist.
2. Runs 100 warm-up iteration then runs inference for 1000 to 2000 iterations for each batch size specified in the script, selecting the profile best for each size.
**Note:** The time measurements do not include the time required to copy inputs to the device and copy outputs to the host.
To run the inference benchmark script, run:
```bash
bash scripts/inference_benchmark.sh --gpu <arch>
```
Options for `<arch>` are: 'Volta', 'Xavier', 'Turing', 'Ampere'
Note: Some of the configurations in the benchmark script require 16GB of GPU memory. On GPUs with smaller amounts of memory, parts of the benchmark may fail to run.
Also note that BERT Large engines, especially using mixed precision with large batch sizes and sequence lengths may take a couple hours to build.
### Results
The following sections provide details on how we achieved our performance and inference.
#### Inference performance: NVIDIA A100 (40GB)
Our results were obtained by running the `scripts/inference_benchmark.sh --gpu Ampere` script in the container generated by the TensorRT OSS Dockerfile on NVIDIA A100 with (1x A100 40G) GPUs.
##### BERT Base
| Sequence Length | Batch Size | INT8 Latency (ms) | | | FP16 Latency (ms) | | |
|-----------------|------------|-----------------|-----------------|---------|-----------------|-----------------|---------|
| | | 95th Percentile | 99th Percentile | Average | 95th Percentile | 99th Percentile | Average |
| 128 | 1 | 0.77 | 0.77 | 0.77 | 0.78 | 0.80 | 0.78 |
| 128 | 2 | 0.76 | 0.77 | 0.76 | 0.92 | 0.93 | 0.92 |
| 128 | 4 | 0.93 | 1.18 | 0.93 | 1.19 | 1.51 | 1.19 |
| 128 | 8 | 1.19 | 1.20 | 1.19 | 1.78 | 1.78 | 1.77 |
| 128 | 12 | 1.57 | 1.57 | 1.56 | 2.07 | 2.08 | 2.05 |
| 128 | 16 | 1.88 | 1.89 | 1.88 | 2.54 | 2.60 | 2.52 |
| 128 | 24 | 2.65 | 2.65 | 2.64 | 3.65 | 3.70 | 3.61 |
| 128 | 32 | 3.21 | 3.22 | 3.21 | 4.71 | 4.74 | 4.67 |
| 128 | 64 | 5.69 | 5.70 | 5.64 | 8.87 | 8.96 | 8.81 |
| 128 | 128 | 10.84 | 10.85 | 10.70 | 17.61 | 17.62 | 17.44 |
| 384 | 1 | 1.34 | 1.35 | 1.34 | 1.46 | 1.46 | 1.45 |
| 384 | 2 | 1.56 | 1.79 | 1.56 | 1.85 | 1.85 | 1.84 |
| 384 | 4 | 2.02 | 2.03 | 2.02 | 2.46 | 2.46 | 2.45 |
| 384 | 8 | 2.94 | 2.95 | 2.94 | 3.91 | 3.92 | 3.86 |
| 384 | 12 | 4.07 | 4.07 | 4.06 | 5.54 | 5.55 | 5.47 |
| 384 | 16 | 5.22 | 5.23 | 5.21 | 7.78 | 7.79 | 7.69 |
| 384 | 24 | 7.42 | 7.42 | 7.37 | 10.75 | 10.76 | 10.63 |
| 384 | 32 | 9.92 | 9.93 | 9.77 | 14.58 | 14.73 | 14.52 |
| 384 | 64 | 18.74 | 18.78 | 18.61 | 28.66 | 28.70 | 28.39 |
| 384 | 128 | 36.40 | 36.42 | 36.05 | 55.36 | 55.90 | 55.21 |
##### BERT Large
| Sequence Length | Batch Size | INT8 Latency (ms) | | | FP16 Latency (ms) | | |
|-----------------|------------|-----------------|-----------------|---------|-----------------|-----------------|---------|
| | | 95th Percentile | 99th Percentile | Average | 95th Percentile | 99th Percentile | Average |
| 128 | 1 | 1.60 | 1.61 | 1.60 | 1.87 | 1.88 | 1.87 |
| 128 | 2 | 1.94 | 1.95 | 1.94 | 2.36 | 2.37 | 2.35 |
| 128 | 4 | 2.45 | 2.46 | 2.44 | 3.36 | 3.37 | 3.36 |
| 128 | 8 | 3.82 | 3.83 | 3.79 | 4.98 | 5.00 | 4.95 |
| 128 | 12 | 4.22 | 4.23 | 4.22 | 6.45 | 6.46 | 6.38 |
| 128 | 16 | 5.75 | 5.75 | 5.74 | 8.50 | 8.53 | 8.43 |
| 128 | 24 | 7.10 | 7.11 | 7.04 | 11.47 | 11.49 | 11.31 |
| 128 | 32 | 9.61 | 9.61 | 9.51 | 15.49 | 15.50 | 15.25 |
| 128 | 64 | 17.25 | 17.25 | 17.11 | 29.43 | 29.73 | 29.29 |
| 128 | 128 | 33.25 | 33.58 | 33.05 | 56.98 | 57.17 | 56.68 |
| 384 | 1 | 3.00 | 3.01 | 2.99 | 3.52 | 3.53 | 3.51 |
| 384 | 2 | 3.71 | 3.72 | 3.71 | 4.97 | 4.99 | 4.97 |
| 384 | 4 | 5.08 | 5.09 | 5.08 | 7.01 | 7.01 | 6.92 |
| 384 | 8 | 9.04 | 9.05 | 9.04 | 12.71 | 12.72 | 12.67 |
| 384 | 12 | 11.65 | 11.71 | 11.57 | 18.24 | 18.25 | 18.04 |
| 384 | 16 | 15.63 | 15.63 | 15.49 | 24.24 | 24.28 | 23.94 |
| 384 | 24 | 22.57 | 22.61 | 22.36 | 35.77 | 35.78 | 35.38 |
| 384 | 32 | 29.66 | 29.66 | 29.33 | 47.09 | 47.11 | 46.81 |
| 384 | 64 | 57.20 | 57.34 | 56.93 | 92.12 | 92.49 | 91.61 |
| 384 | 128 | 112.00 | 112.42 | 111.24 | 180.61 | 181.02 | 179.56 |
#### Inference performance: NVIDIA T4 (16GB)
Our results were obtained by running the `scripts/inference_benchmark.sh --gpu Turing` script in the container generated by the TensorRT OSS Dockerfile on NVIDIA T4 with (1x T4 16G) GPUs.
##### BERT Base
| Sequence Length | Batch Size | INT8 Latency (ms) | | | FP16 Latency (ms) | | |
|-----------------|------------|-----------------|-----------------|---------|-----------------|-----------------|---------|
| | | 95th Percentile | 99th Percentile | Average | 95th Percentile | 99th Percentile | Average |
| 128 | 1 | 1.67 | 1.67 | 1.66 | 1.82 | 1.96 | 1.76 |
| 128 | 2 | 1.94 | 1.95 | 1.89 | 2.58 | 2.67 | 2.50 |
| 128 | 4 | 2.73 | 2.80 | 2.64 | 4.30 | 4.34 | 4.17 |
| 128 | 8 | 4.93 | 4.96 | 4.81 | 8.85 | 9.74 | 8.36 |
| 128 | 12 | 6.85 | 7.05 | 6.70 | 12.83 | 13.19 | 12.34 |
| 128 | 16 | 9.65 | 9.89 | 9.43 | 17.70 | 18.27 | 17.01 |
| 128 | 24 | 15.04 | 15.70 | 14.68 | 27.00 | 27.87 | 26.50 |
| 128 | 32 | 20.55 | 21.01 | 19.88 | 34.51 | 34.81 | 33.83 |
| 128 | 64 | 40.48 | 41.29 | 39.87 | 67.84 | 68.57 | 67.03 |
| 128 | 128 | 82.17 | 82.53 | 80.95 | 132.78 | 133.23 | 131.64 |
| 384 | 1 | 2.75 | 2.78 | 2.67 | 3.73 | 3.79 | 3.63 |
| 384 | 2 | 4.22 | 4.38 | 4.09 | 6.68 | 7.27 | 6.53 |
| 384 | 4 | 7.87 | 8.07 | 7.75 | 13.22 | 13.50 | 12.83 |
| 384 | 8 | 16.07 | 16.13 | 15.77 | 28.01 | 28.72 | 27.48 |
| 384 | 12 | 23.87 | 24.15 | 23.53 | 40.96 | 41.51 | 39.39 |
| 384 | 16 | 31.87 | 32.25 | 30.99 | 51.56 | 51.83 | 51.00 |
| 384 | 24 | 48.14 | 48.33 | 47.22 | 82.06 | 82.56 | 80.13 |
| 384 | 32 | 64.07 | 64.48 | 63.19 | 102.64 | 103.33 | 101.20 |
| 384 | 64 | 129.58 | 130.37 | 125.79 | 215.79 | 216.38 | 213.87 |
| 384 | 128 | 258.69 | 259.74 | 245.91 | 414.96 | 415.57 | 413.16 |
##### BERT Large
| Sequence Length | Batch Size | INT8 Latency (ms) | | | FP16 Latency (ms) | | |
|-----------------|------------|-----------------|-----------------|---------|-----------------|-----------------|---------|
| | | 95th Percentile | 99th Percentile | Average | 95th Percentile | 99th Percentile | Average |
| 128 | 1 | 4.20 | 4.35 | 4.10 | 5.05 | 5.21 | 4.91 |
| 128 | 2 | 5.41 | 5.70 | 5.30 | 7.99 | 8.31 | 7.79 |
| 128 | 4 | 8.48 | 8.68 | 8.32 | 14.87 | 15.28 | 14.44 |
| 128 | 8 | 15.20 | 15.22 | 14.91 | 29.66 | 30.20 | 28.97 |
| 128 | 12 | 23.54 | 23.72 | 23.21 | 45.48 | 45.90 | 44.91 |
| 128 | 16 | 31.04 | 31.38 | 30.46 | 62.06 | 62.61 | 60.27 |
| 128 | 24 | 48.00 | 48.59 | 47.44 | 84.17 | 84.50 | 83.43 |
| 128 | 32 | 64.41 | 64.77 | 63.54 | 113.60 | 113.98 | 112.32 |
| 128 | 64 | 128.03 | 128.45 | 126.36 | 223.89 | 224.83 | 220.75 |
| 128 | 128 | 246.96 | 247.80 | 245.00 | 441.52 | 442.26 | 439.65 |
| 384 | 1 | 7.88 | 8.06 | 7.73 | 11.84 | 12.11 | 11.51 |
| 384 | 2 | 13.00 | 13.18 | 12.80 | 23.59 | 24.13 | 23.12 |
| 384 | 4 | 25.14 | 25.32 | 24.70 | 46.66 | 46.69 | 45.81 |
| 384 | 8 | 50.14 | 50.65 | 49.41 | 86.74 | 87.47 | 85.40 |
| 384 | 12 | 72.92 | 73.01 | 71.86 | 127.10 | 127.44 | 125.66 |
| 384 | 16 | 97.00 | 97.26 | 95.47 | 169.41 | 169.93 | 167.55 |
| 384 | 24 | 149.70 | 150.28 | 148.00 | 258.26 | 258.88 | 255.79 |
| 384 | 32 | 192.74 | 193.85 | 190.59 | 339.87 | 340.55 | 337.86 |
| 384 | 64 | 385.85 | 387.66 | 383.62 | 692.10 | 692.88 | 689.73 |
| 384 | 128 | 780.95 | 781.81 | 778.82 | 1367.61 | 1368.85 | 1365.16 |
#### Inference performance: NVIDIA V100 (16GB)
Our results were obtained by running the `scripts/inference_benchmark.sh --gpu Volta` script in the container generated by the TensorRT OSS Dockerfile on NVIDIA V100 with (1x V100 16G) GPUs.
##### BERT Base
| Sequence Length | Batch Size | INT8 Latency (ms) | | | FP16 Latency (ms) | | |
|-----------------|------------|-----------------|-----------------|---------|-----------------|-----------------|---------|
| | | 95th Percentile | 99th Percentile | Average | 95th Percentile | 99th Percentile | Average |
| 128 | 1 | 1.39 | 1.39 | 1.39 | 1.23 | 1.23 | 1.23 |
| 128 | 2 | 1.76 | 1.76 | 1.75 | 1.49 | 1.49 | 1.48 |
| 128 | 4 | 2.35 | 2.36 | 2.34 | 2.12 | 2.13 | 2.11 |
| 128 | 8 | 3.69 | 3.7 | 3.65 | 3.35 | 3.36 | 3.32 |
| 128 | 12 | 4.79 | 4.83 | 4.75 | 4.65 | 4.67 | 4.61 |
| 128 | 16 | 6.7 | 6.72 | 6.64 | 6.3 | 6.35 | 6.25 |
| 128 | 24 | 8.95 | 8.96 | 8.9 | 8.68 | 8.71 | 8.6 |
| 128 | 32 | 14.74 | 14.77 | 14.59 | 14.16 | 14.18 | 14.06 |
| 128 | 64 | 24.12 | 24.14 | 23.98 | 22.57 | 22.63 | 22.47 |
| 128 | 128 | 45.59 | 45.65 | 45.53 | 43.45 | 43.51 | 43.25 |
| 384 | 1 | 2.17 | 2.18 | 2.16 | 1.98 | 1.98 | 1.97 |
| 384 | 2 | 3.4 | 3.42 | 3.38 | 3.11 | 3.11 | 3.08 |
| 384 | 4 | 5.61 | 5.62 | 5.57 | 5.5 | 5.52 | 5.46 |
| 384 | 8 | 10.58 | 10.63 | 10.49 | 10.26 | 10.29 | 10.17 |
| 384 | 12 | 16.55 | 16.57 | 16.43 | 15.8 | 15.83 | 15.69 |
| 384 | 16 | 21.15 | 21.19 | 21.04 | 20.09 | 20.12 | 19.94 |
| 384 | 24 | 30.95 | 31 | 30.77 | 29.44 | 29.51 | 29.24 |
| 384 | 32 | 47.94 | 48.03 | 47.66 | 47.97 | 48.05 | 47.56 |
| 384 | 64 | 81.8 | 81.91 | 81.62 | 76.84 | 77.05 | 76.4 |
| 384 | 128 | 159.87 | 160.06 | 159.47 | 151.4 | 151.61 | 150.85 |
##### BERT Large
| Sequence Length | Batch Size | INT8 Latency (ms) | | | FP16 Latency (ms) | | |
|-----------------|------------|-----------------|-----------------|---------|-----------------|-----------------|---------|
| | | 95th Percentile | 99th Percentile | Average | 95th Percentile | 99th Percentile | Average |
| 128 | 1 | 3.43 | 3.44 | 3.42 | 3.06 | 3.07 | 3.05 |
| 128 | 2 | 4.35 | 4.37 | 4.33 | 3.79 | 3.8 | 3.79 |
| 128 | 4 | 6.8 | 6.83 | 6.74 | 6.02 | 6.05 | 5.98 |
| 128 | 8 | 11 | 11.07 | 10.93 | 10.57 | 10.62 | 10.47 |
| 128 | 12 | 16.28 | 16.31 | 16.15 | 15.06 | 15.1 | 14.96 |
| 128 | 16 | 20.33 | 20.44 | 20.13 | 20.47 | 20.51 | 20.25 |
| 128 | 24 | 30.63 | 30.66 | 30.33 | 28.65 | 28.8 | 28.48 |
| 128 | 32 | 45.28 | 45.35 | 45.09 | 46.88 | 47.02 | 46.43 |
| 128 | 64 | 75.33 | 75.57 | 74.82 | 71.88 | 71.97 | 71.47 |
| 128 | 128 | 148.1 | 148.31 | 147.59 | 140.81 | 140.97 | 140.35 |
| 384 | 1 | 6.16 | 6.17 | 6.12 | 5.7 | 5.72 | 5.66 |
| 384 | 2 | 10.25 | 10.27 | 10.18 | 9.46 | 9.49 | 9.37 |
| 384 | 4 | 18.44 | 18.5 | 18.27 | 17.22 | 17.28 | 17.09 |
| 384 | 8 | 34.67 | 34.71 | 34.41 | 32.71 | 32.79 | 32.45 |
| 384 | 12 | 49.04 | 49.13 | 48.79 | 47.53 | 47.77 | 47.27 |
| 384 | 16 | 67.08 | 67.21 | 66.75 | 62.86 | 63.01 | 62.76 |
| 384 | 24 | 94.22 | 94.39 | 94.04 | 92.08 | 92.2 | 91.86 |
| 384 | 32 | 148.96 | 149.11 | 148.59 | 147.7 | 147.84 | 147.23 |
| 384 | 64 | 245.91 | 246.09 | 244.67 | 240.16 | 240.43 | 239.07 |
+678
View File
@@ -0,0 +1,678 @@
#!/usr/bin/env python3
#
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
import argparse
import ctypes
import json
import numpy as np
import os
import os.path
import re
import sys
import time
import onnx
# TensorRT
import tensorrt as trt
from helpers.calibrator import BertCalibrator as BertCalibrator
try:
from tensorflow.python import pywrap_tensorflow as pyTF
except ImportError as err:
sys.stderr.write("""Error: Failed to import tensorflow module ({})\n""".format(err))
sys.exit()
"""
TensorRT Initialization
"""
TRT_LOGGER = trt.Logger(trt.Logger.INFO)
handle = ctypes.CDLL("libnvinfer_plugin.so", mode=ctypes.RTLD_GLOBAL)
if not handle:
raise RuntimeError("Could not load plugin library. Is `libnvinfer_plugin.so` on your LD_LIBRARY_PATH?")
trt.init_libnvinfer_plugins(TRT_LOGGER, "")
plg_registry = trt.get_plugin_registry()
emln_plg_creator = plg_registry.get_plugin_creator("CustomEmbLayerNormPluginDynamic", "1", "")
qkv2_plg_creator = plg_registry.get_plugin_creator("CustomQKVToContextPluginDynamic", "1", "")
skln_plg_creator = plg_registry.get_plugin_creator("CustomSkipLayerNormPluginDynamic", "1", "")
fc_plg_creator = plg_registry.get_plugin_creator("CustomFCPluginDynamic", "1", "")
"""
Attentions Keys
"""
WQ = "self_query_kernel"
BQ = "self_query_bias"
WK = "self_key_kernel"
BK = "self_key_bias"
WV = "self_value_kernel"
BV = "self_value_bias"
WQKV = "self_qkv_kernel"
BQKV = "self_qkv_bias"
"""
Transformer Keys
"""
W_AOUT = "attention_output_dense_kernel"
B_AOUT = "attention_output_dense_bias"
AOUT_LN_BETA = "attention_output_layernorm_beta"
AOUT_LN_GAMMA = "attention_output_layernorm_gamma"
W_MID = "intermediate_dense_kernel"
B_MID = "intermediate_dense_bias"
W_LOUT = "output_dense_kernel"
B_LOUT = "output_dense_bias"
LOUT_LN_BETA = "output_layernorm_beta"
LOUT_LN_GAMMA = "output_layernorm_gamma"
"""
Squad Output Keys
"""
SQD_W = "squad_output_weights"
SQD_B = "squad_output_bias"
class BertConfig:
def __init__(self, bert_config_path, use_fp16, use_int8, use_strict, use_fc2_gemm, use_int8_skipln, use_int8_multihead, use_qat):
with open(bert_config_path, "r") as f:
data = json.load(f)
self.num_attention_heads = data["num_attention_heads"]
self.hidden_size = data["hidden_size"]
self.intermediate_size = data["intermediate_size"]
self.num_hidden_layers = data["num_hidden_layers"]
self.head_size = self.hidden_size // self.num_attention_heads
self.use_fp16 = use_fp16
self.use_int8 = use_int8
self.use_fc2_gemm = use_fc2_gemm
self.use_strict = use_strict
self.use_int8_skipln = use_int8_skipln
self.use_int8_multihead = use_int8_multihead
self.is_calib_mode = False
self.use_qat = use_qat
def set_tensor_name(tensor, prefix, name):
tensor.name = prefix + name
def set_output_name(layer, prefix, name, out_idx = 0):
set_tensor_name(layer.get_output(out_idx), prefix, name)
def set_output_range(layer, maxval, out_idx = 0):
layer.get_output(out_idx).set_dynamic_range(-maxval, maxval)
def attention_layer_opt(prefix, config, init_dict, network, input_tensor, imask):
"""
Add the attention layer
"""
assert(len(input_tensor.shape) == 5)
B, S, hidden_size, _, _ = input_tensor.shape
num_heads = config.num_attention_heads
head_size = int(hidden_size / num_heads)
Wall = init_dict[prefix + WQKV]
Ball = init_dict[prefix + BQKV]
# FC_attention
if config.use_int8:
mult_all = network.add_convolution(input_tensor, 3 * hidden_size, (1, 1), Wall, Ball)
else:
mult_all = network.add_fully_connected(input_tensor, 3 * hidden_size, Wall, Ball)
if config.use_qat:
dr_qkv = max(
init_dict[prefix + 'self_qv_a_input_quantizer_amax'],
init_dict[prefix + 'self_qv_b_input_quantizer_amax'],
init_dict[prefix + 'self_av_b_input_quantizer_amax'],
)
set_output_range(mult_all, dr_qkv)
set_output_name(mult_all, prefix, "qkv_mult")
has_mask = imask is not None
# QKV2CTX
dtype = trt.float32
if config.use_fp16:
dtype = trt.float16
# Multi-head attention doesn't use INT8 inputs and output by default unless it is specified.
if config.use_int8 and config.use_int8_multihead and not config.is_calib_mode:
dtype = trt.int8
pf_type = trt.PluginField("type_id", np.array([int(dtype)], np.int32), trt.PluginFieldType.INT32)
pf_hidden_size = trt.PluginField("hidden_size", np.array([hidden_size], np.int32), trt.PluginFieldType.INT32)
pf_num_heads = trt.PluginField("num_heads", np.array([num_heads], np.int32), trt.PluginFieldType.INT32)
pf_has_mask = trt.PluginField("has_mask", np.array([has_mask], np.int32), trt.PluginFieldType.INT32)
if config.use_qat:
dr_probs = init_dict[prefix + 'self_av_a_input_quantizer_amax']
dq_probs = dr_probs / 127.0
pf_dq_probs = trt.PluginField("dq_probs", np.array([dq_probs], np.float32), trt.PluginFieldType.FLOAT32)
pfc = trt.PluginFieldCollection([pf_hidden_size, pf_num_heads, pf_has_mask, pf_type, pf_dq_probs])
else:
pfc = trt.PluginFieldCollection([pf_hidden_size, pf_num_heads, pf_has_mask, pf_type])
qkv2ctx_plug = qkv2_plg_creator.create_plugin("qkv2ctx", pfc)
qkv_in = [mult_all.get_output(0)]
if has_mask:
qkv_in.append(imask)
qkv2ctx = network.add_plugin_v2(qkv_in, qkv2ctx_plug)
if config.use_qat:
dr_ctx = init_dict[prefix + 'output_dense_input_amax']
set_output_range(qkv2ctx, dr_ctx)
set_output_name(qkv2ctx, prefix, "context_layer")
return qkv2ctx
def skipln(prefix, config, init_dict, network, input_tensor, skip, bias=None):
"""
Add the skip layer
"""
idims = input_tensor.shape
assert len(idims) == 5
hidden_size = idims[2]
dtype = trt.float32
if config.use_fp16:
dtype = trt.float16
# Skip layernorm doesn't use INT8 inputs and output by default unless it is specified.
if config.use_int8 and config.use_int8_skipln and not config.is_calib_mode:
dtype = trt.int8
pf_ld = trt.PluginField("ld", np.array([hidden_size], np.int32), trt.PluginFieldType.INT32)
wbeta = init_dict[prefix + "beta"]
pf_beta = trt.PluginField("beta", wbeta.numpy(), trt.PluginFieldType.FLOAT32)
wgamma = init_dict[prefix + "gamma"]
pf_gamma = trt.PluginField("gamma", wgamma.numpy(), trt.PluginFieldType.FLOAT32)
pf_type = trt.PluginField("type_id", np.array([int(dtype)], np.int32), trt.PluginFieldType.INT32)
fields = [pf_ld, pf_beta, pf_gamma, pf_type ]
if bias:
pf_bias = trt.PluginField("bias", bias.numpy(), trt.PluginFieldType.FLOAT32)
fields.append(pf_bias)
pfc = trt.PluginFieldCollection(fields)
skipln_plug = skln_plg_creator.create_plugin("skipln", pfc)
skipln_inputs = [input_tensor, skip]
layer = network.add_plugin_v2(skipln_inputs, skipln_plug)
return layer
def custom_fc(config, network, input_tensor, out_dims, W):
pf_out_dims = trt.PluginField("out_dims", np.array([out_dims], dtype=np.int32), trt.PluginFieldType.INT32)
pf_W = trt.PluginField("W", W.numpy(), trt.PluginFieldType.FLOAT32)
pf_type = trt.PluginField("type_id", np.array([1 if config.use_fp16 else 0], np.int32), trt.PluginFieldType.INT32)
pfc = trt.PluginFieldCollection([pf_out_dims, pf_W, pf_type])
fc_plugin = fc_plg_creator.create_plugin("fcplugin", pfc)
plug_inputs = [input_tensor]
out_dense = network.add_plugin_v2(plug_inputs, fc_plugin)
return out_dense
def transformer_layer_opt(prefix, config, init_dict, network, input_tensor, imask):
"""
Add the transformer layer
"""
idims = input_tensor.shape
assert len(idims) == 5
hidden_size = idims[2]
if config.use_qat:
dr_input = init_dict[prefix + 'attention_self_query_input_amax']
assert(dr_input ==init_dict[prefix + 'attention_self_key_input_amax'] )
assert(dr_input ==init_dict[prefix + 'attention_self_value_input_amax'] )
input_tensor.set_dynamic_range(-dr_input, dr_input)
context_transposed = attention_layer_opt(prefix + "attention_", config, init_dict, network, input_tensor, imask)
attention_heads = context_transposed.get_output(0)
# FC0
B_aout = init_dict[prefix + B_AOUT]
if config.use_int8:
W_aout = init_dict[prefix + W_AOUT]
attention_out_fc = network.add_convolution(attention_heads, hidden_size, (1, 1), W_aout, B_aout)
B_aout = None
if not config.use_int8_skipln:
attention_out_fc.set_output_type(0, trt.DataType.HALF if config.use_fp16 else trt.DataType.FLOAT)
if config.use_qat:
dr_fc_aout = init_dict[prefix + 'attention_output_add_local_input_quantizer_amax']
set_output_range(attention_out_fc, dr_fc_aout)
else:
W_aoutT = init_dict[prefix + W_AOUT + "_notrans"]
attention_out_fc = custom_fc(config, network, attention_heads, hidden_size, W_aoutT)
skiplayer = skipln(prefix + "attention_output_layernorm_",config, init_dict, network, attention_out_fc.get_output(0), input_tensor, B_aout)
attention_ln = skiplayer.get_output(0)
if config.use_qat:
dr_skln1 = init_dict[prefix + 'intermediate_dense_input_amax']
set_output_range(skiplayer, dr_skln1)
# FC1 + GELU
B_mid = init_dict[prefix + B_MID]
W_mid = init_dict[prefix + W_MID]
if config.use_int8:
mid_dense = network.add_convolution(attention_ln, config.intermediate_size, (1, 1), W_mid, B_mid)
else:
mid_dense = network.add_fully_connected(attention_ln, config.intermediate_size, W_mid, B_mid)
mid_dense_out = mid_dense.get_output(0)
POW = network.add_constant((1, 1, 1, 1, 1), trt.Weights(np.ascontiguousarray([3.0], dtype=np.float32)))
MULTIPLY = network.add_constant((1, 1, 1, 1, 1), trt.Weights(np.ascontiguousarray([0.044715], dtype=np.float32)))
SQRT = network.add_constant((1, 1, 1, 1, 1), trt.Weights((np.ascontiguousarray([0.79788456080286535587989211986876], dtype=np.float32))))
ONE = network.add_constant((1, 1, 1, 1, 1), trt.Weights((np.ascontiguousarray([1.0], dtype=np.float32))))
HALF = network.add_constant((1, 1, 1, 1, 1), trt.Weights((np.ascontiguousarray([0.5], dtype=np.float32))))
X_pow = network.add_elementwise(mid_dense_out, POW.get_output(0), trt.ElementWiseOperation.POW)
X_pow_t = X_pow.get_output(0)
X_mul = network.add_elementwise(X_pow_t, MULTIPLY.get_output(0), trt.ElementWiseOperation.PROD)
X_add = network.add_elementwise(mid_dense_out, X_mul.get_output(0), trt.ElementWiseOperation.SUM)
X_sqrt = network.add_elementwise(X_add.get_output(0), SQRT.get_output(0), trt.ElementWiseOperation.PROD)
X_sqrt_tensor = X_sqrt.get_output(0)
X_tanh = network.add_activation(X_sqrt_tensor, trt.ActivationType.TANH)
X_tanh_tensor = X_tanh.get_output(0)
X_one = network.add_elementwise(X_tanh_tensor, ONE.get_output(0), trt.ElementWiseOperation.SUM)
CDF = network.add_elementwise(X_one.get_output(0), HALF.get_output(0), trt.ElementWiseOperation.PROD)
gelu_layer = network.add_elementwise(CDF.get_output(0), mid_dense_out, trt.ElementWiseOperation.PROD)
intermediate_act = gelu_layer.get_output(0)
set_tensor_name(intermediate_act, prefix, "gelu")
if config.use_int8:
if config.use_qat:
dr_gelu = init_dict[prefix + 'output_dense_input_amax']
set_output_range(gelu_layer, dr_gelu)
else:
# use gelu10 according to whitepaper http://arxiv.org/abs/2004.09602
set_output_range(gelu_layer, 10)
# FC2
# Dense to hidden size
B_lout = init_dict[prefix + B_LOUT]
if config.use_int8 and not config.use_fc2_gemm:
W_lout = init_dict[prefix + W_LOUT]
out_dense = network.add_convolution(intermediate_act, hidden_size, (1, 1), W_lout, B_lout)
B_lout = None
if not config.use_int8_skipln:
out_dense.set_output_type(0, trt.DataType.HALF if config.use_fp16 else trt.DataType.FLOAT)
else:
W_loutT = init_dict[prefix + W_LOUT + "_notrans"]
out_dense = custom_fc(config, network, intermediate_act, hidden_size, W_loutT)
if config.use_qat:
dr_fc_out = init_dict[prefix + 'output_add_local_input_quantizer_amax']
set_output_range(out_dense, dr_fc_out)
set_output_name(out_dense, prefix + "output_", "dense")
out_layer = skipln(prefix + "output_layernorm_", config, init_dict, network, out_dense.get_output(0), attention_ln, B_lout)
set_output_name(out_layer, prefix + "output_", "reshape")
return out_layer
def bert_model(config, init_dict, network, input_tensor, input_mask):
"""
Create the bert model
"""
prev_input = input_tensor
for layer in range(0, config.num_hidden_layers):
ss = "l{}_".format(layer)
out_layer = transformer_layer_opt(ss, config, init_dict, network, prev_input, input_mask)
prev_input = out_layer.get_output(0)
if config.use_qat:
dr_out = init_dict["bert_encoder_final_input_quantizer_amax"]
set_output_range(out_layer, dr_out)
return prev_input
def squad_output(prefix, config, init_dict, network, input_tensor):
"""
Create the squad output
"""
idims = input_tensor.shape
assert len(idims) == 5
B, S, hidden_size, _, _ = idims
W_out = init_dict[prefix + SQD_W]
B_out = init_dict[prefix + SQD_B]
W = network.add_constant((1, hidden_size, 2), W_out)
dense = network.add_fully_connected(input_tensor, 2, W_out, B_out)
OUT = network.add_shuffle(dense.get_output(0))
OUT.second_transpose = (1, 0, 2, 3, 4)
set_output_name(OUT, prefix, "squad_logits")
return OUT
def load_tf_weights(inputbase, config):
"""
Load the weights from the tensorflow checkpoint
"""
weights_dict = dict()
try:
reader = pyTF.NewCheckpointReader(inputbase)
tensor_dict = reader.get_variable_to_shape_map()
# There might be training-related variables in the checkpoint that can be discarded
param_names = [key for key in sorted(tensor_dict) if "adam" not in key and "global_step" not in key and "pooler" not in key]
count = len(param_names)
TRT_LOGGER.log(TRT_LOGGER.INFO, "Found {:} entries in weight map".format(count))
for pn in param_names:
toks = pn.lower().split("/")
if "encoder" in pn:
assert ("layer" in pn)
l = (re.findall("\d+", pn))[0]
outname = "l{}_".format(l) + "_".join(toks[3:])
else:
outname = "_".join(toks)
tensor = reader.get_tensor(pn)
shape = tensor.shape
if pn.find("kernel") != -1:
weights_dict[outname + "_notrans"] = trt.Weights(np.ascontiguousarray(tensor).flatten())
TRT_LOGGER.log(TRT_LOGGER.VERBOSE, "Transposing {}\n".format(np))
tensor = np.transpose(tensor)
shape = tensor.shape
flat_tensor = tensor.flatten()
shape_str = "{} ".format(len(shape)) + " ".join([str(d) for d in shape])
weights_dict[outname] = trt.Weights(flat_tensor)
TRT_LOGGER.log(TRT_LOGGER.VERBOSE, "Original name: {:}, TensorRT name: {:}, shape: {:}".format(pn, outname, shape_str))
N = config.num_attention_heads
H = config.head_size
additional_dict = dict()
for key, value in weights_dict.items():
pos = key.find(BQ)
if pos != -1:
hidden_size = value.size
prefix = key[:pos]
Bq_ = value
Bk_ = weights_dict[prefix + BK]
Bv_ = weights_dict[prefix + BV]
Wq_ = weights_dict[prefix + WQ]
Wk_ = weights_dict[prefix + WK]
Wv_ = weights_dict[prefix + WV]
mat_size = hidden_size * hidden_size
wcount = 3 * mat_size
Wall = np.zeros(wcount, np.float32)
bcount = 3 * hidden_size
Ball = np.zeros(bcount, np.float32)
Wall[0:mat_size] = Wq_.numpy()[0:mat_size]
Wall[mat_size:2*mat_size] = Wk_.numpy()[0:mat_size]
Wall[2*mat_size:3*mat_size] = Wv_.numpy()[0:mat_size]
Ball[0:hidden_size] = Bq_.numpy()[0:hidden_size]
Ball[hidden_size:2*hidden_size] = Bk_.numpy()[0:hidden_size]
Ball[2*hidden_size:3*hidden_size] = Bv_.numpy()[0:hidden_size]
Wall = np.ascontiguousarray(Wall.reshape((3, N, H, N, H)).transpose((1, 0, 2, 3, 4)), dtype=np.float32)
Ball = np.ascontiguousarray(Ball.reshape((3, N, H)).transpose((1, 0, 2)), dtype=np.float32)
additional_dict[prefix + WQKV] = trt.Weights(Wall)
additional_dict[prefix + BQKV] = trt.Weights(Ball)
additional_dict[prefix + WQKV + "_notrans"] = trt.Weights(Wall.T)
except Exception as error:
TRT_LOGGER.log(TRT_LOGGER.ERROR, str(error))
weights_dict.update(additional_dict)
return weights_dict
def onnx_to_trt_name(onnx_name):
"""
Converting variables in the onnx checkpoint to names corresponding to the naming convention used in the TF version, expected by the builder
"""
onnx_name = onnx_name.lower()
toks = [t.strip('_') for t in onnx_name.split('.')]
if toks[0] == 'bert': #embeddings or encoder
if toks[1] == 'encoder': #transformer
if toks[-2] == 'layernorm': #bias->beta, weight->gamma
toks[-1] = 'beta' if toks[-1] == 'bias' else 'gamma'
elif (toks[-2] == 'dense' or toks[-2] in {'key', 'value', 'query'}) and toks[-1] == 'weight':
toks[-1] = 'kernel'
elif (toks[-3] == 'dense' or toks[-3] in {'key', 'value', 'query'}) and toks[-1] == 'amax':
if toks[-2] == 'weight_quantizer':
toks[-2] = 'kernel'
elif toks[-2] == 'input_quantizer':
toks[-2] = 'input'
if 'final_input_quantizer' not in toks[2]:
toks = toks[3:]
toks[0] = 'l{}'.format(int(toks[0]))
else:
if toks[-2] == 'layernorm': #bias->beta, weight->gamma
toks[-1] = 'beta' if toks[-1] == 'bias' else 'gamma'
else: #embeddings: drop "_weight" suffix
if toks[-1] == 'amax':
toks[-2] = 'amax'
toks = toks[:-1]
elif 'qa' in onnx_name:
name = 'cls_squad_output_bias' if toks[-1] == 'bias' else 'cls_squad_output_weights'
return name
else:
print("Encountered unknown case:", onnx_name)
assert(False)
parsed = '_'.join(toks)
return parsed
def load_onnx_weights_and_quant(path, config):
"""
Load the weights from the onnx checkpoint
"""
N = config.num_attention_heads
H = config.head_size
hidden_size = config.hidden_size
model = onnx.load(path)
weights = model.graph.initializer
tensor_dict = dict([(onnx_to_trt_name(w.name), np.frombuffer(w.raw_data, np.float32).reshape(w.dims)) for w in weights])
weights_dict = dict()
for outname, tensor in tensor_dict.items():
if outname.find("_amax") != -1:
weights_dict[outname] = tensor
elif outname.find(BQ) != -1:
prefix = outname[:outname.find(BQ)]
Wqkv = np.zeros((3, hidden_size, hidden_size), np.float32)
Bqkv = np.zeros((3, hidden_size), np.float32)
Wqkv[0,:,:] = tensor_dict[prefix + WQ]
Wqkv[1,:,:] = tensor_dict[prefix + WK]
Wqkv[2,:,:] = tensor_dict[prefix + WV]
Bqkv[0,:] = tensor
Bqkv[1,:] = tensor_dict[prefix + BK]
Bqkv[2,:] = tensor_dict[prefix + BV]
Wqkv = np.ascontiguousarray(Wqkv.reshape((3, N, H, N, H)).transpose((1,0,2,3,4)))
Bqkv = np.ascontiguousarray(Bqkv.reshape((3, N, H)).transpose((1,0,2)))
weights_dict[prefix + WQKV] = trt.Weights(Wqkv)
weights_dict[prefix + BQKV] = trt.Weights(Bqkv)
weights_dict[prefix + WQKV + "_notrans"] = trt.Weights(Wqkv.T)
elif outname.find(BK) != -1 or outname.find(BV) != -1 or outname.find(WQ) != -1 or outname.find(WK) != -1 or outname.find(WV) != -1:
pass
else:
flat_tensor = np.ascontiguousarray(tensor).flatten()
weights_dict[outname] = trt.Weights(flat_tensor)
if outname.find("kernel") != -1:
tensor = np.transpose(tensor)
weights_dict[outname + "_notrans"] = trt.Weights(np.ascontiguousarray(tensor).flatten())
TRT_LOGGER.log(TRT_LOGGER.INFO, "Found {:} entries in weight map".format(len(weights_dict)))
return weights_dict
def emb_layernorm(builder, network, config, weights_dict, builder_config, sequence_length, batch_sizes):
if len(batch_sizes) > 1:
input_ids = network.add_input(name="input_ids", dtype=trt.int32, shape=(sequence_length, -1))
segment_ids = network.add_input(name="segment_ids", dtype=trt.int32, shape=(sequence_length, -1))
input_mask = network.add_input(name="input_mask", dtype=trt.int32, shape=(sequence_length, -1))
# Specify profiles for the batch sizes we're interested in.
# Make sure the profile also works for all sizes not covered by the previous profile.
prev_size = 0
for batch_size in sorted(batch_sizes):
profile = builder.create_optimization_profile()
min_shape = (sequence_length, prev_size + 1)
shape = (sequence_length, batch_size)
profile.set_shape("input_ids", min=min_shape, opt=shape, max=shape)
profile.set_shape("segment_ids", min=min_shape, opt=shape, max=shape)
profile.set_shape("input_mask", min=min_shape, opt=shape, max=shape)
builder_config.add_optimization_profile(profile)
prev_size = batch_size
else:
input_ids = network.add_input(name="input_ids", dtype=trt.int32, shape=(sequence_length, batch_sizes[0]))
segment_ids = network.add_input(name="segment_ids", dtype=trt.int32, shape=(sequence_length, batch_sizes[0]))
input_mask = network.add_input(name="input_mask", dtype=trt.int32, shape=(sequence_length, batch_sizes[0]))
wbeta = trt.PluginField("bert_embeddings_layernorm_beta", weights_dict["bert_embeddings_layernorm_beta"].numpy(), trt.PluginFieldType.FLOAT32)
wgamma = trt.PluginField("bert_embeddings_layernorm_gamma", weights_dict["bert_embeddings_layernorm_gamma"].numpy(), trt.PluginFieldType.FLOAT32)
wwordemb = trt.PluginField("bert_embeddings_word_embeddings", weights_dict["bert_embeddings_word_embeddings"].numpy(), trt.PluginFieldType.FLOAT32)
wtokemb = trt.PluginField("bert_embeddings_token_type_embeddings", weights_dict["bert_embeddings_token_type_embeddings"].numpy(), trt.PluginFieldType.FLOAT32)
wposemb = trt.PluginField("bert_embeddings_position_embeddings", weights_dict["bert_embeddings_position_embeddings"].numpy(), trt.PluginFieldType.FLOAT32)
output_fp16 = trt.PluginField("output_fp16", np.array([1 if config.use_fp16 else 0]).astype(np.int32), trt.PluginFieldType.INT32)
full_mask = trt.PluginField("full_mask", np.array([1 if config.use_fp16 else 0]).astype(np.int32), trt.PluginFieldType.INT32)
pfc = trt.PluginFieldCollection([wbeta, wgamma, wwordemb, wtokemb, wposemb, output_fp16, full_mask])
fn = emln_plg_creator.create_plugin("embeddings", pfc)
inputs = [input_ids, segment_ids, input_mask]
emb_layer = network.add_plugin_v2(inputs, fn)
if config.use_qat:
set_output_range(emb_layer, 1, 1)
set_output_name(emb_layer, "embeddings_", "output")
return emb_layer
def build_engine(batch_sizes, workspace_size, sequence_length, config, weights_dict, squad_json, vocab_file, calibrationCacheFile, calib_num):
explicit_batch_flag = 1 << int(trt.NetworkDefinitionCreationFlag.EXPLICIT_BATCH)
with trt.Builder(TRT_LOGGER) as builder, builder.create_network(explicit_batch_flag) as network, builder.create_builder_config() as builder_config:
builder_config.max_workspace_size = workspace_size * (1024 * 1024)
if config.use_fp16:
builder_config.set_flag(trt.BuilderFlag.FP16)
if config.use_int8:
builder_config.set_flag(trt.BuilderFlag.INT8)
if not config.use_qat:
calibrator = BertCalibrator(squad_json, vocab_file, calibrationCacheFile, 1, sequence_length, calib_num)
builder_config.set_quantization_flag(trt.QuantizationFlag.CALIBRATE_BEFORE_FUSION)
builder_config.int8_calibrator = calibrator
if config.use_strict:
builder_config.set_flag(trt.BuilderFlag.STRICT_TYPES)
# Create the network
emb_layer = emb_layernorm(builder, network, config, weights_dict, builder_config, sequence_length, batch_sizes)
embeddings = emb_layer.get_output(0)
mask_idx = emb_layer.get_output(1)
bert_out = bert_model(config, weights_dict, network, embeddings, mask_idx)
squad_logits = squad_output("cls_", config, weights_dict, network, bert_out)
squad_logits_out = squad_logits.get_output(0)
network.mark_output(squad_logits_out)
build_start_time = time.time()
engine = builder.build_engine(network, builder_config)
build_time_elapsed = (time.time() - build_start_time)
TRT_LOGGER.log(TRT_LOGGER.INFO, "build engine in {:.3f} Sec".format(build_time_elapsed))
if config.use_int8 and not config.use_qat:
calibrator.free()
return engine
def generate_calibration_cache(sequence_length, workspace_size, config, weights_dict, squad_json, vocab_file, calibrationCacheFile, calib_num):
"""
BERT demo needs a separate engine building path to generate calibration cache.
This is because we need to configure SLN and MHA plugins in FP32 mode when
generating calibration cache, and INT8 mode when building the actual engine.
This cache could be generated by examining certain training data and can be
reused across different configurations.
"""
# dynamic shape not working with calibration, so we need generate a calibration cache first using fulldims network
if not config.use_int8 or os.path.exists(calibrationCacheFile):
return calibrationCacheFile
# generate calibration cache
saved_use_fp16 = config.use_fp16
config.use_fp16 = False
config.is_calib_mode = True
with build_engine([1], workspace_size, sequence_length, config, weights_dict, squad_json, vocab_file, calibrationCacheFile, calib_num) as engine:
TRT_LOGGER.log(TRT_LOGGER.INFO, "calibration cache generated in {:}".format(calibrationCacheFile))
config.use_fp16 = saved_use_fp16
config.is_calib_mode = False
def main():
parser = argparse.ArgumentParser(description="TensorRT BERT Sample", formatter_class=argparse.ArgumentDefaultsHelpFormatter)
parser.add_argument("-m", "--ckpt", required=False,
help="The checkpoint file basename, e.g.: basename(model.ckpt-766908.data-00000-of-00001) is model.ckpt-766908")
parser.add_argument("-x", "--onnx", required=False, help="The ONNX model file path.")
parser.add_argument("-o", "--output", required=True, default="bert_base_384.engine", help="The bert engine file, ex bert.engine")
parser.add_argument("-b", "--batch-size", default=[], action="append", help="Batch size(s) to optimize for. The engine will be usable with any batch size below this, but may not be optimal for smaller sizes. Can be specified multiple times to optimize for more than one batch size.", type=int)
parser.add_argument("-s", "--sequence-length", default=128, help="Sequence length of the BERT model", type=int)
parser.add_argument("-c", "--config-dir", required=True,
help="The folder containing the bert_config.json, which can be downloaded e.g. from https://github.com/google-research/bert#pre-trained-models or by running download_models.py in dle/TensorFlow/LanguageModeling/BERT/data/pretrained_models_google")
parser.add_argument("-f", "--fp16", action="store_true", help="Indicates that inference should be run in FP16 precision", required=False)
parser.add_argument("-i", "--int8", action="store_true", help="Indicates that inference should be run in INT8 precision", required=False)
parser.add_argument("-t", "--strict", action="store_true", help="Indicates that inference should be run in strict precision mode", required=False)
parser.add_argument("-w", "--workspace-size", default=1000, help="Workspace size in MiB for building the BERT engine", type=int)
parser.add_argument("-j", "--squad-json", default="squad/dev-v1.1.json", help="squad json dataset used for int8 calibration", required=False)
parser.add_argument("-v", "--vocab-file", default="./pre-trained_model/uncased_L-24_H-1024_A-16/vocab.txt", help="Path to file containing entire understandable vocab", required=False)
parser.add_argument("-n", "--calib-num", default=100, help="calibration batch numbers", type=int)
parser.add_argument("-p", "--calib-path", help="calibration cache path", required=False)
parser.add_argument("-g", "--force-fc2-gemm", action="store_true", help="Force use gemm to implement FC2 layer", required=False)
parser.add_argument("-iln", "--force-int8-skipln", action="store_true", help="Run skip layernorm with INT8 (FP32 or FP16 by default) inputs and output", required=False)
parser.add_argument("-imh", "--force-int8-multihead", action="store_true", help="Run multi-head attention with INT8 (FP32 or FP16 by default) input and output", required=False)
args, _ = parser.parse_known_args()
args.batch_size = args.batch_size or [1]
bert_config_path = os.path.join(args.config_dir, "bert_config.json")
TRT_LOGGER.log(TRT_LOGGER.INFO, "Using configuration file: {:}".format(bert_config_path))
config = BertConfig(bert_config_path, args.fp16, args.int8, args.strict, args.force_fc2_gemm, args.force_int8_skipln, args.force_int8_multihead, args.int8 and args.onnx != None)
if args.calib_path != None:
calib_cache = args.calib_path
else:
calib_cache = "BertSquadL{}H{}A{}S{}CalibCache".format(config.num_hidden_layers, config.head_size, config.num_attention_heads, args.sequence_length)
if args.onnx != None:
weights_dict = load_onnx_weights_and_quant(args.onnx, config)
elif args.ckpt != None:
weights_dict = load_tf_weights(args.ckpt, config)
generate_calibration_cache(args.sequence_length, args.workspace_size, config, weights_dict, args.squad_json, args.vocab_file, calib_cache, args.calib_num)
else:
raise RuntimeError("You need either specify TF checkpoint using option --ckpt or ONNX using option --onnx to build TRT BERT model.")
with build_engine(args.batch_size, args.workspace_size, args.sequence_length, config, weights_dict, args.squad_json, args.vocab_file, calib_cache, args.calib_num) as engine:
TRT_LOGGER.log(TRT_LOGGER.VERBOSE, "Serializing Engine...")
serialized_engine = engine.serialize()
TRT_LOGGER.log(TRT_LOGGER.INFO, "Saving Engine to {:}".format(args.output))
with open(args.output, "wb") as fout:
fout.write(serialized_engine)
TRT_LOGGER.log(TRT_LOGGER.INFO, "Done.")
if __name__ == "__main__":
main()
+111
View File
@@ -0,0 +1,111 @@
#!/usr/bin/env python3
#
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
import tensorrt as trt
import os
import pycuda.driver as cuda
import pycuda.autoinit
import numpy as np
import helpers.tokenization as tokenization
import helpers.data_processing as dp
class BertCalibrator(trt.IInt8LegacyCalibrator):
def __init__(self, squad_json, vocab_file, cache_file, batch_size, max_seq_length, num_inputs):
# Whenever you specify a custom constructor for a TensorRT class,
# you MUST call the constructor of the parent explicitly.
trt.IInt8LegacyCalibrator.__init__(self)
self.cache_file = cache_file
# Every time get_batch is called, the next batch of size batch_size will be copied to the device and returned.
self.data = dp.read_squad_json(squad_json)
self.max_seq_length = max_seq_length
self.batch_size = batch_size
self.current_index = 0
self.num_inputs = num_inputs
self.tokenizer = tokenization.BertTokenizer(vocab_file=vocab_file, do_lower_case=True)
self.doc_stride = 128
self.max_query_length = 64
# Allocate enough memory for a whole batch.
self.device_inputs = [cuda.mem_alloc(self.max_seq_length * trt.int32.itemsize * self.batch_size) for binding in range(3)]
def free(self):
for dinput in self.device_inputs:
dinput.free()
def get_batch_size(self):
return self.batch_size
# TensorRT passes along the names of the engine bindings to the get_batch function.
# You don't necessarily have to use them, but they can be useful to understand the order of
# the inputs. The bindings list is expected to have the same ordering as 'names'.
def get_batch(self, names):
if self.current_index + self.batch_size > self.num_inputs:
print("Calibrating index {:} batch size {:} exceed max input limit {:} sentences".format(self.current_index, self.batch_size, self.num_inputs))
return None
current_batch = int(self.current_index / self.batch_size)
if current_batch % 10 == 0:
print("Calibrating batch {:}, containing {:} sentences".format(current_batch, self.batch_size))
input_ids = []
segment_ids = []
input_mask = []
for i in range(self.batch_size):
example = self.data[self.current_index + i]
features = dp.convert_example_to_features(example.doc_tokens, example.question_text, self.tokenizer, self.max_seq_length, self.doc_stride, self.max_query_length)
if len(input_ids) and len(segment_ids) and len(input_mask):
input_ids = np.concatenate((input_ids, features[0].input_ids))
segment_ids = np.concatenate((segment_ids, features[0].segment_ids))
input_mask = np.concatenate((input_mask, features[0].input_mask))
else:
input_ids = features[0].input_ids
segment_ids = features[0].segment_ids
input_mask = features[0].input_mask
cuda.memcpy_htod(self.device_inputs[0], input_ids.ravel())
cuda.memcpy_htod(self.device_inputs[1], segment_ids.ravel())
cuda.memcpy_htod(self.device_inputs[2], input_mask.ravel())
self.current_index += self.batch_size
return self.device_inputs
def read_calibration_cache(self):
# If there is a cache, use it instead of calibrating again. Otherwise, implicitly return None.
if os.path.exists(self.cache_file):
with open(self.cache_file, "rb") as f:
return f.read()
def write_calibration_cache(self, cache):
with open(self.cache_file, "wb") as f:
f.write(cache)
f.flush()
os.fsync(f)
def get_quantile(self):
return 0.9999
def get_regression_cutoff(self):
return 1.0
def read_histogram_cache(self, length):
return None
def write_histogram_cache(self, ptr, length):
return None
+495
View File
@@ -0,0 +1,495 @@
#!/usr/bin/env python3
#
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
import helpers.tokenization as tokenization
import collections
import numpy as np
import six
import math
import json
def convert_doc_tokens(paragraph_text):
""" Return the list of tokens from the doc text """
def is_whitespace(c):
if c == " " or c == "\t" or c == "\r" or c == "\n" or ord(c) == 0x202F:
return True
return False
doc_tokens = []
prev_is_whitespace = True
for c in paragraph_text:
if is_whitespace(c):
prev_is_whitespace = True
else:
if prev_is_whitespace:
doc_tokens.append(c)
else:
doc_tokens[-1] += c
prev_is_whitespace = False
return doc_tokens
def _check_is_max_context(doc_spans, cur_span_index, position):
"""Check if this is the 'max context' doc span for the token."""
# Because of the sliding window approach taken to scoring documents, a single
# token can appear in multiple documents. E.g.
# Doc: the man went to the store and bought a gallon of milk
# Span A: the man went to the
# Span B: to the store and bought
# Span C: and bought a gallon of
# ...
#
# Now the word 'bought' will have two scores from spans B and C. We only
# want to consider the score with "maximum context", which we define as
# the *minimum* of its left and right context (the *sum* of left and
# right context will always be the same, of course).
#
# In the example the maximum context for 'bought' would be span C since
# it has 1 left context and 3 right context, while span B has 4 left context
# and 0 right context.
best_score = None
best_span_index = None
for (span_index, doc_span) in enumerate(doc_spans):
end = doc_span.start + doc_span.length - 1
if position < doc_span.start:
continue
if position > end:
continue
num_left_context = position - doc_span.start
num_right_context = end - position
score = min(num_left_context, num_right_context) + 0.01 * doc_span.length
if best_score is None or score > best_score:
best_score = score
best_span_index = span_index
return cur_span_index == best_span_index
def convert_example_to_features(doc_tokens, question_text, tokenizer, max_seq_length,
doc_stride, max_query_length):
"""Loads a data file into a list of `InputBatch`s."""
query_tokens = tokenizer.tokenize(question_text)
if len(query_tokens) > max_query_length:
query_tokens = query_tokens[0:max_query_length]
tok_to_orig_index = []
orig_to_tok_index = []
all_doc_tokens = []
for (i, token) in enumerate(doc_tokens):
orig_to_tok_index.append(len(all_doc_tokens))
sub_tokens = tokenizer.tokenize(token)
for sub_token in sub_tokens:
tok_to_orig_index.append(i)
all_doc_tokens.append(sub_token)
# The -3 accounts for [CLS], [SEP] and [SEP]
max_tokens_for_doc = max_seq_length - len(query_tokens) - 3
# We can have documents that are longer than the maximum sequence length.
# To deal with this we do a sliding window approach, where we take chunks
# of the up to our max length with a stride of `doc_stride`.
_DocSpan = collections.namedtuple( # pylint: disable=invalid-name
"DocSpan", ["start", "length"])
doc_spans = []
start_offset = 0
while start_offset < len(all_doc_tokens):
length = len(all_doc_tokens) - start_offset
if length > max_tokens_for_doc:
length = max_tokens_for_doc
doc_spans.append(_DocSpan(start=start_offset, length=length))
if start_offset + length == len(all_doc_tokens):
break
start_offset += min(length, doc_stride)
_Feature = collections.namedtuple( # pylint: disable=invalid-name
"Feature",
["input_ids", "input_mask", "segment_ids", "tokens", "token_to_orig_map", "token_is_max_context"])
features = []
for (doc_span_index, doc_span) in enumerate(doc_spans):
tokens = []
token_to_orig_map = {}
token_is_max_context = {}
segment_ids = []
tokens.append("[CLS]")
segment_ids.append(0)
for token in query_tokens:
tokens.append(token)
segment_ids.append(0)
tokens.append("[SEP]")
segment_ids.append(0)
for i in range(doc_span.length):
split_token_index = doc_span.start + i
token_to_orig_map[len(tokens)] = tok_to_orig_index[split_token_index]
is_max_context = _check_is_max_context(doc_spans, doc_span_index, split_token_index)
token_is_max_context[len(tokens)] = is_max_context
tokens.append(all_doc_tokens[split_token_index])
segment_ids.append(1)
tokens.append("[SEP]")
segment_ids.append(1)
input_ids = tokenizer.convert_tokens_to_ids(tokens)
# The mask has 1 for real tokens and 0 for padding tokens. Only real
# tokens are attended to.
input_mask = [1] * len(input_ids)
# Zero-pad up to the sequence length.
while len(input_ids) < max_seq_length:
input_ids.append(0)
input_mask.append(0)
segment_ids.append(0)
assert len(input_ids) == max_seq_length
assert len(input_mask) == max_seq_length
assert len(segment_ids) == max_seq_length
def create_int_feature(values):
feature = np.asarray(values, dtype=np.int32, order=None)
return feature
features.append(_Feature(
input_ids = create_int_feature(input_ids),
input_mask = create_int_feature(input_mask),
segment_ids = create_int_feature(segment_ids),
tokens = tokens,
token_to_orig_map = token_to_orig_map,
token_is_max_context = token_is_max_context
))
return features
def read_squad_json(input_file):
"""read from squad json into a list of examples"""
with open(input_file, "r", encoding='utf-8') as reader:
input_data = json.load(reader)["data"]
_Example = collections.namedtuple( # pylint: disable=invalid-name
"Example",
["id", "question_text", "doc_tokens"])
examples = []
for entry in input_data:
for paragraph in entry["paragraphs"]:
paragraph_text = paragraph["context"]
doc_tokens = convert_doc_tokens(paragraph_text)
for qa in paragraph["qas"]:
examples.append(_Example(
id = qa["id"],
question_text = qa["question"],
doc_tokens = doc_tokens
))
return examples
def _get_best_indexes(logits, n_best_size):
"""Get the n-best logits from a list."""
index_and_score = sorted(enumerate(logits), key=lambda x: x[1], reverse=True)
best_indexes = []
for i in range(len(index_and_score)):
if i >= n_best_size:
break
best_indexes.append(index_and_score[i][0])
return best_indexes
def get_final_text(pred_text, orig_text, do_lower_case):
"""Project the tokenized prediction back to the original text."""
# When we created the data, we kept track of the alignment between original
# (whitespace tokenized) tokens and our WordPiece tokenized tokens. So
# now `orig_text` contains the span of our original text corresponding to the
# span that we predicted.
#
# However, `orig_text` may contain extra characters that we don't want in
# our prediction.
#
# For example, let's say:
# pred_text = steve smith
# orig_text = Steve Smith's
#
# We don't want to return `orig_text` because it contains the extra "'s".
#
# We don't want to return `pred_text` because it's already been normalized
# (the SQuAD eval script also does punctuation stripping/lower casing but
# our tokenizer does additional normalization like stripping accent
# characters).
#
# What we really want to return is "Steve Smith".
#
# Therefore, we have to apply a semi-complicated alignment heruistic between
# `pred_text` and `orig_text` to get a character-to-charcter alignment. This
# can fail in certain cases in which case we just return `orig_text`.
def _strip_spaces(text):
ns_chars = []
ns_to_s_map = collections.OrderedDict()
for (i, c) in enumerate(text):
if c == " ":
continue
ns_to_s_map[len(ns_chars)] = i
ns_chars.append(c)
ns_text = "".join(ns_chars)
return (ns_text, ns_to_s_map)
# We first tokenize `orig_text`, strip whitespace from the result
# and `pred_text`, and check if they are the same length. If they are
# NOT the same length, the heuristic has failed. If they are the same
# length, we assume the characters are one-to-one aligned.
tokenizer = tokenization.BasicTokenizer(do_lower_case=do_lower_case)
tok_text = " ".join(tokenizer.tokenize(orig_text))
start_position = tok_text.find(pred_text)
if start_position == -1:
return orig_text
end_position = start_position + len(pred_text) - 1
(orig_ns_text, orig_ns_to_s_map) = _strip_spaces(orig_text)
(tok_ns_text, tok_ns_to_s_map) = _strip_spaces(tok_text)
if len(orig_ns_text) != len(tok_ns_text):
return orig_text
# We then project the characters in `pred_text` back to `orig_text` using
# the character-to-character alignment.
tok_s_to_ns_map = {}
for (i, tok_index) in six.iteritems(tok_ns_to_s_map):
tok_s_to_ns_map[tok_index] = i
orig_start_position = None
if start_position in tok_s_to_ns_map:
ns_start_position = tok_s_to_ns_map[start_position]
if ns_start_position in orig_ns_to_s_map:
orig_start_position = orig_ns_to_s_map[ns_start_position]
if orig_start_position is None:
return orig_text
orig_end_position = None
if end_position in tok_s_to_ns_map:
ns_end_position = tok_s_to_ns_map[end_position]
if ns_end_position in orig_ns_to_s_map:
orig_end_position = orig_ns_to_s_map[ns_end_position]
if orig_end_position is None:
return orig_text
output_text = orig_text[orig_start_position:(orig_end_position + 1)]
return output_text
def _compute_softmax(scores):
"""Compute softmax probability over raw logits."""
if not scores:
return []
max_score = None
for score in scores:
if max_score is None or score > max_score:
max_score = score
exp_scores = []
total_sum = 0.0
for score in scores:
x = math.exp(score - max_score)
exp_scores.append(x)
total_sum += x
probs = []
for score in exp_scores:
probs.append(score / total_sum)
return probs
def get_predictions(doc_tokens, features, results, n_best_size, max_answer_length):
_PrelimPrediction = collections.namedtuple( # pylint: disable=invalid-name
"PrelimPrediction",
["feature_index", "start_index", "end_index", "start_logit", "end_logit"])
prediction = ""
scores_diff_json = 0.0
prelim_predictions = []
# keep track of the minimum score of null start+end of position 0
score_null = 1000000 # large and positive
min_null_feature_index = 0 # the paragraph slice with min mull score
null_start_logit = 0 # the start logit at the slice with min null score
null_end_logit = 0 # the end logit at the slice with min null score
version_2_with_negative = False
for result in results:
start_indexes = _get_best_indexes(result.start_logits, n_best_size)
end_indexes = _get_best_indexes(result.end_logits, n_best_size)
feature = features[result.feature_index]
# if we could have irrelevant answers, get the min score of irrelevant
if version_2_with_negative:
feature_null_score = result.start_logits[0] + result.end_logits[0]
if feature_null_score < score_null:
score_null = feature_null_score
min_null_feature_index = 0
null_start_logit = result.start_logits[0]
null_end_logit = result.end_logits[0]
for start_index in start_indexes:
for end_index in end_indexes:
# We could hypothetically create invalid predictions, e.g., predict
# that the start of the span is in the question. We throw out all
# invalid predictions.
if start_index >= len(feature.tokens):
continue
if end_index >= len(feature.tokens):
continue
if start_index not in feature.token_to_orig_map:
continue
if end_index not in feature.token_to_orig_map:
continue
if not feature.token_is_max_context.get(start_index, False):
continue
if end_index < start_index:
continue
length = end_index - start_index + 1
if length > max_answer_length:
continue
prelim_predictions.append(
_PrelimPrediction(
feature_index=result.feature_index,
start_index=start_index,
end_index=end_index,
start_logit=result.start_logits[start_index],
end_logit=result.end_logits[end_index]))
if version_2_with_negative:
prelim_predictions.append(
_PrelimPrediction(
feature_index=result.feature_index,
start_index=0,
end_index=0,
start_logit=null_start_logit,
end_logit=null_end_logit))
prelim_predictions = sorted(
prelim_predictions,
key=lambda x: (x.start_logit + x.end_logit),
reverse=True)
_NbestPrediction = collections.namedtuple( # pylint: disable=invalid-name
"NbestPrediction", ["text", "start_logit", "end_logit"])
seen_predictions = {}
nbest = []
for pred in prelim_predictions:
if len(nbest) >= n_best_size:
break
if pred.start_index > 0: # this is a non-null prediction
feature = features[pred.feature_index]
tok_tokens = feature.tokens[pred.start_index:(pred.end_index + 1)]
orig_doc_start = feature.token_to_orig_map[pred.start_index]
orig_doc_end = feature.token_to_orig_map[pred.end_index]
orig_tokens = doc_tokens[orig_doc_start:(orig_doc_end + 1)]
tok_text = " ".join(tok_tokens)
# De-tokenize WordPieces that have been split off.
tok_text = tok_text.replace(" ##", "")
tok_text = tok_text.replace("##", "")
# Clean whitespace
tok_text = tok_text.strip()
tok_text = " ".join(tok_text.split())
orig_text = " ".join(orig_tokens)
final_text = get_final_text(tok_text, orig_text, True)
if final_text in seen_predictions:
continue
seen_predictions[final_text] = True
else:
final_text = ""
seen_predictions[final_text] = True
if len(final_text):
nbest.append(
_NbestPrediction(
text=final_text,
start_logit=pred.start_logit,
end_logit=pred.end_logit))
# if we didn't inlude the empty option in the n-best, inlcude it
if version_2_with_negative:
if "" not in seen_predictions:
nbest.append(
_NbestPrediction(
text="", start_logit=null_start_logit,
end_logit=null_end_logit))
# In very rare edge cases we could have no valid predictions. So we
# just create a nonce prediction in this case to avoid failure.
if not nbest:
nbest.append(
_NbestPrediction(text="empty", start_logit=0.0, end_logit=0.0))
assert len(nbest) >= 1
total_scores = []
best_non_null_entry = None
for entry in nbest:
total_scores.append(entry.start_logit + entry.end_logit)
if not best_non_null_entry:
if entry.text:
best_non_null_entry = entry
probs = _compute_softmax(total_scores)
nbest_json = []
for (i, entry) in enumerate(nbest):
output = collections.OrderedDict()
output["text"] = entry.text
output["probability"] = probs[i]
output["start_logit"] = entry.start_logit
output["end_logit"] = entry.end_logit
nbest_json.append(output)
assert len(nbest_json) >= 1
null_score_diff_threshold = 0.0
if not version_2_with_negative:
prediction = nbest_json[0]["text"]
else:
# predict "" iff the null score - the score of best non-null > threshold
score_diff = score_null - best_non_null_entry.start_logit - (
best_non_null_entry.end_logit)
scores_diff_json = score_diff
if score_diff > null_score_diff_threshold:
prediction = ""
else:
prediction = best_non_null_entry.text
return prediction, nbest_json, scores_diff_json
+429
View File
@@ -0,0 +1,429 @@
# coding=utf-8
# Copyright 2018 The Google AI Language Team Authors.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""Tokenization classes."""
from __future__ import absolute_import
from __future__ import division
from __future__ import print_function
import collections
import re
import unicodedata
import six
def validate_case_matches_checkpoint(do_lower_case, init_checkpoint):
"""Checks whether the casing config is consistent with the checkpoint name."""
# The casing has to be passed in by the user and there is no explicit check
# as to whether it matches the checkpoint. The casing information probably
# should have been stored in the bert_config.json file, but it's not, so
# we have to heuristically detect it to validate.
if not init_checkpoint:
return
m = re.match("^.*?([A-Za-z0-9_-]+)/bert_model.ckpt", init_checkpoint)
if m is None:
return
model_name = m.group(1)
lower_models = [
"uncased_L-24_H-1024_A-16", "uncased_L-12_H-768_A-12",
"multilingual_L-12_H-768_A-12", "chinese_L-12_H-768_A-12"
]
cased_models = [
"cased_L-12_H-768_A-12", "cased_L-24_H-1024_A-16",
"multi_cased_L-12_H-768_A-12"
]
is_bad_config = False
if model_name in lower_models and not do_lower_case:
is_bad_config = True
actual_flag = "False"
case_name = "lowercased"
opposite_flag = "True"
if model_name in cased_models and do_lower_case:
is_bad_config = True
actual_flag = "True"
case_name = "cased"
opposite_flag = "False"
if is_bad_config:
raise ValueError(
"You passed in `--do_lower_case=%s` with `--init_checkpoint=%s`. "
"However, `%s` seems to be a %s model, so you "
"should pass in `--do_lower_case=%s` so that the fine-tuning matches "
"how the model was pre-training. If this error is wrong, please "
"just comment out this check." % (actual_flag, init_checkpoint,
model_name, case_name, opposite_flag))
def convert_to_unicode(text):
"""Converts `text` to Unicode (if it's not already), assuming utf-8 input."""
if six.PY3:
if isinstance(text, str):
return text
elif isinstance(text, bytes):
return text.decode("utf-8", "ignore")
else:
raise ValueError("Unsupported string type: %s" % (type(text)))
elif six.PY2:
if isinstance(text, str):
return text.decode("utf-8", "ignore")
elif isinstance(text, unicode):
return text
else:
raise ValueError("Unsupported string type: %s" % (type(text)))
else:
raise ValueError("Not running on Python2 or Python 3?")
def printable_text(text):
"""Returns text encoded in a way suitable for print or `tf.logging`."""
# These functions want `str` for both Python2 and Python3, but in one case
# it's a Unicode string and in the other it's a byte string.
if six.PY3:
if isinstance(text, str):
return text
elif isinstance(text, bytes):
return text.decode("utf-8", "ignore")
else:
raise ValueError("Unsupported string type: %s" % (type(text)))
elif six.PY2:
if isinstance(text, str):
return text
elif isinstance(text, unicode):
return text.encode("utf-8")
else:
raise ValueError("Unsupported string type: %s" % (type(text)))
else:
raise ValueError("Not running on Python2 or Python 3?")
def load_vocab(vocab_file):
"""Loads a vocabulary file into a dictionary."""
vocab = collections.OrderedDict()
index = 0
with open(vocab_file, "r", encoding='utf-8') as reader:
while True:
token = convert_to_unicode(reader.readline())
if not token:
break
token = token.strip()
vocab[token] = index
index += 1
return vocab
def convert_by_vocab(vocab, items):
"""Converts a sequence of [tokens|ids] using the vocab."""
output = []
for item in items:
output.append(vocab[item])
return output
def convert_tokens_to_ids(vocab, tokens):
return convert_by_vocab(vocab, tokens)
def convert_ids_to_tokens(inv_vocab, ids):
return convert_by_vocab(inv_vocab, ids)
def whitespace_tokenize(text):
"""Runs basic whitespace cleaning and splitting on a piece of text."""
text = text.strip()
if not text:
return []
tokens = text.split()
return tokens
class FullTokenizer(object):
"""Runs end-to-end tokenziation."""
def __init__(self, vocab_file, do_lower_case=True):
self.vocab = load_vocab(vocab_file)
self.inv_vocab = {v: k for k, v in self.vocab.items()}
self.basic_tokenizer = BasicTokenizer(do_lower_case=do_lower_case)
self.wordpiece_tokenizer = WordpieceTokenizer(vocab=self.vocab)
def tokenize(self, text):
split_tokens = []
for token in self.basic_tokenizer.tokenize(text):
for sub_token in self.wordpiece_tokenizer.tokenize(token):
split_tokens.append(sub_token)
return split_tokens
def convert_tokens_to_ids(self, tokens):
return convert_by_vocab(self.vocab, tokens)
def convert_ids_to_tokens(self, ids):
return convert_by_vocab(self.inv_vocab, ids)
class BertTokenizer(object):
"""Runs end-to-end tokenization: punctuation splitting + wordpiece"""
def __init__(self, vocab_file, do_lower_case=True):
self.vocab = load_vocab(vocab_file)
self.ids_to_tokens = collections.OrderedDict(
[(ids, tok) for tok, ids in self.vocab.items()])
self.basic_tokenizer = BasicTokenizer(do_lower_case=do_lower_case)
self.wordpiece_tokenizer = WordpieceTokenizer(vocab=self.vocab)
def tokenize(self, text):
split_tokens = []
for token in self.basic_tokenizer.tokenize(text):
for sub_token in self.wordpiece_tokenizer.tokenize(token):
split_tokens.append(sub_token)
return split_tokens
def convert_tokens_to_ids(self, tokens):
"""Converts a sequence of tokens into ids using the vocab."""
ids = []
for token in tokens:
ids.append(self.vocab[token])
return ids
def convert_ids_to_tokens(self, ids):
"""Converts a sequence of ids in wordpiece tokens using the vocab."""
tokens = []
for i in ids:
tokens.append(self.ids_to_tokens[i])
return tokens
class BasicTokenizer(object):
"""Runs basic tokenization (punctuation splitting, lower casing, etc.)."""
def __init__(self, do_lower_case=True):
"""Constructs a BasicTokenizer.
Args:
do_lower_case: Whether to lower case the input.
"""
self.do_lower_case = do_lower_case
def tokenize(self, text):
"""Tokenizes a piece of text."""
text = convert_to_unicode(text)
text = self._clean_text(text)
# This was added on November 1st, 2018 for the multilingual and Chinese
# models. This is also applied to the English models now, but it doesn't
# matter since the English models were not trained on any Chinese data
# and generally don't have any Chinese data in them (there are Chinese
# characters in the vocabulary because Wikipedia does have some Chinese
# words in the English Wikipedia.).
text = self._tokenize_chinese_chars(text)
orig_tokens = whitespace_tokenize(text)
split_tokens = []
for token in orig_tokens:
if self.do_lower_case:
token = token.lower()
token = self._run_strip_accents(token)
split_tokens.extend(self._run_split_on_punc(token))
output_tokens = whitespace_tokenize(" ".join(split_tokens))
return output_tokens
def _run_strip_accents(self, text):
"""Strips accents from a piece of text."""
text = unicodedata.normalize("NFD", text)
output = []
for char in text:
cat = unicodedata.category(char)
if cat == "Mn":
continue
output.append(char)
return "".join(output)
def _run_split_on_punc(self, text):
"""Splits punctuation on a piece of text."""
chars = list(text)
i = 0
start_new_word = True
output = []
while i < len(chars):
char = chars[i]
if _is_punctuation(char):
output.append([char])
start_new_word = True
else:
if start_new_word:
output.append([])
start_new_word = False
output[-1].append(char)
i += 1
return ["".join(x) for x in output]
def _tokenize_chinese_chars(self, text):
"""Adds whitespace around any CJK character."""
output = []
for char in text:
cp = ord(char)
if self._is_chinese_char(cp):
output.append(" ")
output.append(char)
output.append(" ")
else:
output.append(char)
return "".join(output)
def _is_chinese_char(self, cp):
"""Checks whether CP is the codepoint of a CJK character."""
# This defines a "chinese character" as anything in the CJK Unicode block:
# https://en.wikipedia.org/wiki/CJK_Unified_Ideographs_(Unicode_block)
#
# Note that the CJK Unicode block is NOT all Japanese and Korean characters,
# despite its name. The modern Korean Hangul alphabet is a different block,
# as is Japanese Hiragana and Katakana. Those alphabets are used to write
# space-separated words, so they are not treated specially and handled
# like the all of the other languages.
if ((cp >= 0x4E00 and cp <= 0x9FFF) or #
(cp >= 0x3400 and cp <= 0x4DBF) or #
(cp >= 0x20000 and cp <= 0x2A6DF) or #
(cp >= 0x2A700 and cp <= 0x2B73F) or #
(cp >= 0x2B740 and cp <= 0x2B81F) or #
(cp >= 0x2B820 and cp <= 0x2CEAF) or
(cp >= 0xF900 and cp <= 0xFAFF) or #
(cp >= 0x2F800 and cp <= 0x2FA1F)): #
return True
return False
def _clean_text(self, text):
"""Performs invalid character removal and whitespace cleanup on text."""
output = []
for char in text:
cp = ord(char)
if cp == 0 or cp == 0xfffd or _is_control(char):
continue
if _is_whitespace(char):
output.append(" ")
else:
output.append(char)
return "".join(output)
class WordpieceTokenizer(object):
"""Runs WordPiece tokenziation."""
def __init__(self, vocab, unk_token="[UNK]", max_input_chars_per_word=200):
self.vocab = vocab
self.unk_token = unk_token
self.max_input_chars_per_word = max_input_chars_per_word
def tokenize(self, text):
"""Tokenizes a piece of text into its word pieces.
This uses a greedy longest-match-first algorithm to perform tokenization
using the given vocabulary.
For example:
input = "unaffable"
output = ["un", "##aff", "##able"]
Args:
text: A single token or whitespace separated tokens. This should have
already been passed through `BasicTokenizer.
Returns:
A list of wordpiece tokens.
"""
text = convert_to_unicode(text)
output_tokens = []
for token in whitespace_tokenize(text):
chars = list(token)
if len(chars) > self.max_input_chars_per_word:
output_tokens.append(self.unk_token)
continue
is_bad = False
start = 0
sub_tokens = []
while start < len(chars):
end = len(chars)
cur_substr = None
while start < end:
substr = "".join(chars[start:end])
if start > 0:
substr = "##" + substr
if substr in self.vocab:
cur_substr = substr
break
end -= 1
if cur_substr is None:
is_bad = True
break
sub_tokens.append(cur_substr)
start = end
if is_bad:
output_tokens.append(self.unk_token)
else:
output_tokens.extend(sub_tokens)
return output_tokens
def _is_whitespace(char):
"""Checks whether `chars` is a whitespace character."""
# \t, \n, and \r are technically contorl characters but we treat them
# as whitespace since they are generally considered as such.
if char == " " or char == "\t" or char == "\n" or char == "\r":
return True
cat = unicodedata.category(char)
if cat == "Zs":
return True
return False
def _is_control(char):
"""Checks whether `chars` is a control character."""
# These are technically control characters but we count them as whitespace
# characters.
if char == "\t" or char == "\n" or char == "\r":
return False
cat = unicodedata.category(char)
if cat.startswith("C"):
return True
return False
def _is_punctuation(char):
"""Checks whether `chars` is a punctuation character."""
cp = ord(char)
# We treat all non-letter/number ASCII as punctuation.
# Characters such as "^", "$", and "`" are not in the Unicode
# Punctuation class but we treat them as punctuation anyways, for
# consistency.
if ((cp >= 33 and cp <= 47) or (cp >= 58 and cp <= 64) or
(cp >= 91 and cp <= 96) or (cp >= 123 and cp <= 126)):
return True
cat = unicodedata.category(char)
if cat.startswith("P"):
return True
return False
+164
View File
@@ -0,0 +1,164 @@
#
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
#
import argparse
import ctypes
import numpy as np
import tensorrt as trt
import pycuda.driver as cuda
import pycuda.autoinit
import numpy as np
import helpers.tokenization as tokenization
import helpers.data_processing as dp
import pdb
TRT_LOGGER = trt.Logger(trt.Logger.ERROR)
class DeviceBuffer(object):
def __init__(self, shape, dtype=trt.int32):
self.buf = cuda.mem_alloc(trt.volume(shape) * dtype.itemsize)
def binding(self):
return int(self.buf)
def free(self):
self.buf.free()
def main():
parser = argparse.ArgumentParser(description='BERT Inference Benchmark')
parser.add_argument("-e", "--engine", help='Path to BERT TensorRT engine')
parser.add_argument('-b', '--batch-size', default=[], action="append", help='Batch size(s) to benchmark. Can be specified multiple times for more than one batch size. This script assumes that the engine has been built with one optimization profile for each batch size, and that these profiles are in order of increasing batch size.', type=int)
parser.add_argument('-s', '--sequence-length', default=128, help='Sequence length of the BERT model', type=int)
parser.add_argument('-i', '--iterations', default=1, help='Number of iterations to run when benchmarking each batch size.', type=int)
parser.add_argument('-w', '--warm-up-runs', default=0, help='Number of iterations to run prior to benchmarking.', type=int)
parser.add_argument('-r', '--random-seed', required=False, default=12345, help='Random seed.', type=int)
parser.add_argument('-p', '--passage', nargs='*', help='Text for paragraph/passage for BERT QA', default='')
parser.add_argument('-q', '--question', nargs='*', help='Text for query/question for BERT QA', default='')
parser.add_argument('-v', '--vocab-file', help='Path to file containing entire understandable vocab')
args, _ = parser.parse_known_args()
args.batch_size = args.batch_size or [1]
# Import necessary plugins for BERT TensorRT
ctypes.CDLL("libnvinfer_plugin.so", mode=ctypes.RTLD_GLOBAL)
with open(args.engine, 'rb') as f, trt.Runtime(TRT_LOGGER) as runtime, runtime.deserialize_cuda_engine(f.read()) as engine, engine.create_execution_context() as context:
# Allocate buffers large enough to store the largest batch size
max_input_shape = (args.sequence_length, max(args.batch_size))
max_output_shape = (args.sequence_length, max(args.batch_size), 2, 1, 1)
buffers = [
DeviceBuffer(max_input_shape),
DeviceBuffer(max_input_shape),
DeviceBuffer(max_input_shape),
DeviceBuffer(max_output_shape)
]
def question_features(tokens, question):
# Extract features from the paragraph and question
tokenizer = tokenization.FullTokenizer(vocab_file=args.vocab_file, do_lower_case=True)
return dp.convert_example_to_features(tokens, question, tokenizer, args.sequence_length, 128, 64)
# Prepare random input
pseudo_vocab_size = 30522
pseudo_type_vocab_size = 2
np.random.seed(args.random_seed)
paragraph_text = ' '.join(args.passage)
question_text = ' '.join(args.question)
print("\nPassage: {}".format(paragraph_text))
print("\nQuestion: {}".format(question_text))
doc_tokens = dp.convert_doc_tokens(paragraph_text)
features = question_features(doc_tokens, question_text)
test_word_ids = features[0].input_ids
test_segment_ids = features[0].segment_ids
test_input_mask = features[0].input_mask
# Copy input h2d
cuda.memcpy_htod(buffers[0].buf, test_word_ids.ravel())
cuda.memcpy_htod(buffers[1].buf, test_segment_ids.ravel())
cuda.memcpy_htod(buffers[2].buf, test_input_mask.ravel())
num_binding_per_profile = engine.num_bindings // engine.num_optimization_profiles
bench_times = {}
for idx, batch_size in enumerate(sorted(args.batch_size)):
context.active_optimization_profile = idx
# Each profile has unique bindings
binding_idx_offset = idx * num_binding_per_profile
bindings = [0] * binding_idx_offset + [buf.binding() for buf in buffers]
shapes = {
"input_ids": (args.sequence_length, batch_size),
"segment_ids": (args.sequence_length, batch_size),
"input_mask": (args.sequence_length, batch_size),
}
for binding, shape in shapes.items():
context.set_binding_shape(engine[binding] + binding_idx_offset, shape)
assert context.all_binding_shapes_specified
# Inference
total_time = 0
start = cuda.Event()
end = cuda.Event()
stream = cuda.Stream()
# Warmup
for _ in range(args.warm_up_runs):
context.execute_async_v2(bindings=bindings, stream_handle=stream.handle)
stream.synchronize()
# Timing loop
times = []
for _ in range(args.iterations):
start.record(stream)
context.execute_async_v2(bindings=bindings, stream_handle=stream.handle)
end.record(stream)
stream.synchronize()
times.append(end.time_since(start))
# Transfer predictions back from GPU
cuda.memcpy_dtoh_async(h_output, d_output, stream)
for index, batch in enumerate(h_output):
# Data Post-processing
networkOutputs.append(_NetworkOutput(
start_logits = np.array(batch.squeeze()[:, 0]),
end_logits = np.array(batch.squeeze()[:, 1]),
feature_index = feature_index
))
# Compute average time, 95th percentile time and 99th percentile time.
bench_times[batch_size] = times
[b.free() for b in buffers]
for batch_size, times in bench_times.items():
total_time = sum(times)
avg_time = total_time / float(len(times))
times.sort()
percentile95 = times[int(len(times) * 0.95)]
percentile99 = times[int(len(times) * 0.99)]
print("Running {:} iterations with Batch Size: {:}\n\tTotal Time: {:} ms \tAverage Time: {:} ms\t95th Percentile Time: {:} ms\t99th Percentile Time: {:}".format(args.iterations, batch_size, total_time, avg_time, percentile95, percentile99))
if __name__ == '__main__':
main()
+295
View File
@@ -0,0 +1,295 @@
/*
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef INFER_C_BERT_INFER_H
#define INFER_C_BERT_INFER_H
#include "common.h"
#include "logging.h"
#include <NvInfer.h>
#include <NvInferPlugin.h>
#include <algorithm>
#include <cuda_runtime.h>
#include <fstream>
#include <numeric>
#include <vector>
using namespace nvinfer1;
struct BertInference
{
BertInference(
const std::string& enginePath, const int maxBatchSize, const int seqLength, const bool enableGraph = false)
: mSeqLength(seqLength)
, mEnableGraph(enableGraph)
{
gLogInfo << "--------------------\n";
gLogInfo << "Using BERT inference C++\n";
if (enableGraph)
{
gLogInfo << "CUDA Graph is enabled\n";
}
else
{
gLogInfo << "CUDA Graph is disabled\n";
}
gLogInfo << "--------------------\n";
initLibNvInferPlugins(&gLogger, "");
gLogInfo << "Loading BERT Inference Engine ... \n";
std::ifstream input(enginePath, std::ios::binary);
if (!input)
{
gLogError << "Error opening engine file: " << enginePath << "\n";
exit(-1);
}
input.seekg(0, input.end);
const size_t fsize = input.tellg();
input.seekg(0, input.beg);
std::vector<char> bytes(fsize);
input.read(bytes.data(), fsize);
auto runtime = TrtUniquePtr<IRuntime>(createInferRuntime(gLogger));
if (runtime == nullptr)
{
gLogError << "Error creating TRT runtime\n";
exit(-1);
}
mEngine = TrtUniquePtr<ICudaEngine>(runtime->deserializeCudaEngine(bytes.data(), bytes.size(), nullptr));
if (mEngine == nullptr)
{
gLogError << "Error deserializing CUDA engine\n";
exit(-1);
}
gLogInfo << "Done\n";
mContext = TrtUniquePtr<IExecutionContext>(mEngine->createExecutionContext());
if (!mContext)
{
gLogError << "Error creating execution context\n";
exit(-1);
}
gpuErrChk(cudaStreamCreate(&mStream));
allocateBindings(maxBatchSize);
}
void allocateBindings(const int maxBatchSize)
{
const size_t allocationSize = mSeqLength * maxBatchSize * sizeof(int32_t);
// Static sizes with implicit batch size: allocation sizes known to engine
for (int i = 0; i < kBERT_INPUT_NUM; i++)
{
void* devBuf;
gpuErrChk(cudaMalloc(&devBuf, allocationSize));
gpuErrChk(cudaMemset(devBuf, 0, allocationSize));
mDeviceBuffers.emplace_back(devBuf);
mInputSizes.emplace_back(allocationSize);
}
const size_t numOutputItems = maxBatchSize * mSeqLength * 2;
mOutputSize = numOutputItems * sizeof(float);
mOutputDims = {maxBatchSize, mSeqLength, 2, 1, 1};
void* devBuf;
gpuErrChk(cudaMalloc(&devBuf, mOutputSize));
gpuErrChk(cudaMemset(devBuf, 0, mOutputSize));
mDeviceBuffers.emplace_back(devBuf);
mHostOutput.resize(numOutputItems);
mBindings.resize(mEngine->getNbBindings());
}
void prepare(int profIdx, int batchSize)
{
mContext->setOptimizationProfile(profIdx);
const int numBindingPerProfile = mEngine->getNbBindings() / mEngine->getNbOptimizationProfiles();
const int bindingIdxOffset = profIdx * numBindingPerProfile;
std::copy(mDeviceBuffers.begin(), mDeviceBuffers.end(), mBindings.begin() + bindingIdxOffset);
for (int i = 0; i < kBERT_INPUT_NUM; i++)
{
mContext->setBindingDimensions(i + bindingIdxOffset, Dims2(mSeqLength, batchSize));
}
if (!mContext->allInputDimensionsSpecified())
{
gLogError << "Not all input dimensions are specified for the exeuction context\n";
exit(-1);
}
if (mEnableGraph)
{
cudaGraph_t graph;
cudaGraphExec_t exec;
// warm up and let mContext do cublas initialization
bool status = mContext->enqueueV2(mBindings.data(), mStream, nullptr);
if (!status)
{
gLogError << "Enqueue failed\n";
exit(-1);
}
gLogVerbose << "Capturing graph\n";
gpuErrChk(cudaStreamBeginCapture(mStream, cudaStreamCaptureModeRelaxed));
status = mContext->enqueueV2(mBindings.data(), mStream, nullptr);
if (!status)
{
gLogError << "Enqueue failed\n";
exit(-1);
}
gpuErrChk(cudaStreamEndCapture(mStream, &graph));
gpuErrChk(cudaStreamSynchronize(mStream));
gpuErrChk(cudaGraphInstantiate(&exec, graph, NULL, NULL, 0));
mExecGraph = exec;
}
}
void run(const void* inputIds, const void* segmentIds, const void* inputMask, int warmUps, int iterations)
{
const std::vector<const void*> inputBuffers = {inputIds, segmentIds, inputMask};
for (int i = 0; i < kBERT_INPUT_NUM; i++)
{
gpuErrChk(
cudaMemcpyAsync(mDeviceBuffers[i], inputBuffers[i], mInputSizes[i], cudaMemcpyHostToDevice, mStream));
}
gLogInfo << "Warming up " << warmUps << " iterations ...\n";
for (int it = 0; it < warmUps; it++)
{
if (mEnableGraph)
{
gpuErrChk(cudaGraphLaunch(mExecGraph, mStream));
}
else
{
bool status = mContext->enqueueV2(mBindings.data(), mStream, nullptr);
if (!status)
{
gLogError << "Enqueue failed\n";
exit(-1);
}
}
}
gpuErrChk(cudaStreamSynchronize(mStream));
cudaEvent_t start, stop;
gpuErrChk(cudaEventCreate(&start));
gpuErrChk(cudaEventCreate(&stop));
std::vector<float> times;
gLogInfo << "Running " << iterations << " iterations ...\n";
for (int it = 0; it < iterations; it++)
{
gpuErrChk(cudaEventRecord(start, mStream));
if (mEnableGraph)
{
gpuErrChk(cudaGraphLaunch(mExecGraph, mStream));
}
else
{
bool status = mContext->enqueueV2(mBindings.data(), mStream, nullptr);
if (!status)
{
gLogError << "Enqueue failed\n";
exit(-1);
}
}
gpuErrChk(cudaEventRecord(stop, mStream));
gpuErrChk(cudaStreamSynchronize(mStream));
float time;
gpuErrChk(cudaEventElapsedTime(&time, start, stop));
times.push_back(time);
}
gpuErrChk(cudaMemcpyAsync(
mHostOutput.data(), mDeviceBuffers[kBERT_INPUT_NUM], mOutputSize, cudaMemcpyDeviceToHost, mStream));
gpuErrChk(cudaStreamSynchronize(mStream));
mTimes.push_back(times);
}
void run(int profIdx, int batchSize, const void* inputIds, const void* segmentIds, const void* inputMask,
int warmUps, int iterations)
{
prepare(profIdx, batchSize);
run(inputIds, segmentIds, inputMask, warmUps, iterations);
}
void reportTiming(int batchIndex, int batchSize)
{
std::vector<float>& times = mTimes[batchIndex];
const float totalTime = std::accumulate(times.begin(), times.end(), 0.0);
const float avgTime = totalTime / times.size();
sort(times.begin(), times.end());
const float percentile95 = times[(int) ((float) times.size() * 0.95)];
const float percentile99 = times[(int) ((float) times.size() * 0.99)];
const int throughput = (int) ((float) batchSize * (1000.0 / avgTime));
gLogInfo << "Running " << times.size() << " iterations with Batch Size: " << batchSize << "\n";
gLogInfo << "\tTotal Time: " << totalTime << " ms \n";
gLogInfo << "\tAverage Time: " << avgTime << " ms\n";
gLogInfo << "\t95th Percentile Time: " << percentile95 << " ms\n";
gLogInfo << "\t99th Percentile Time: " << percentile99 << " ms\n";
gLogInfo << "\tThroughtput: " << throughput << " sentences/s\n";
}
~BertInference()
{
gpuErrChk(cudaStreamDestroy(mStream));
for (auto& buf : mDeviceBuffers)
{
gpuErrChk(cudaFree(buf));
}
}
static const int kBERT_INPUT_NUM = 3;
const int mSeqLength;
const bool mEnableGraph;
TrtUniquePtr<ICudaEngine> mEngine{nullptr};
TrtUniquePtr<IExecutionContext> mContext{nullptr};
std::vector<void*> mBindings;
cudaStream_t mStream{NULL};
std::vector<void*> mDeviceBuffers;
std::vector<float> mHostOutput;
std::vector<size_t> mInputSizes;
size_t mOutputSize;
std::vector<int> mOutputDims;
std::vector<std::vector<float>> mTimes;
cudaGraphExec_t mExecGraph;
};
#endif // INFER_C_BERT_INFER_H
+178
View File
@@ -0,0 +1,178 @@
/*
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef INFER_C_COMMON_H
#define INFER_C_COMMON_H
#include "logging.h"
#include <cuda_runtime_api.h>
#include <getopt.h>
#include <memory>
#include <vector>
struct Args
{
bool help{false};
std::string engine{};
std::vector<int> batchSize;
int sequenceLength{128};
int iterations{200};
int warmUpRuns{10};
int randomSeed{12345};
bool enableGraph{false};
};
//!
//! \brief Populates the Args struct with the provided command-line parameters.
//!
//! \throw invalid_argument if any of the arguments are not valid
//!
//! \return boolean If return value is true, execution can continue, otherwise program should exit
//!
inline bool parseArgs(Args& args, int argc, char* argv[])
{
while (1)
{
int arg;
// clang-format off
static struct option long_options[] =
{
{"help", no_argument, 0, 'h'},
{"engine", required_argument, 0, 'e'},
{"batch_size", required_argument, 0, 'b'},
{"sequence_length", no_argument, 0, 's'},
{"iterations", required_argument, 0, 'i'},
{"warm_up_runs", required_argument, 0, 'w'},
{"ramdon_seed", required_argument, 0, 'r'},
{"enable_graph", no_argument, 0, 'g'},
{nullptr, 0, nullptr, 0}
};
// clang-format on
int option_index = 0;
arg = getopt_long(argc, argv, "he:b:s:i:w:r:g", long_options, &option_index);
if (arg == -1)
{
break;
}
switch (arg)
{
case 'h': args.help = true; return false;
case 'e':
if (optarg)
{
args.engine = optarg;
}
else
{
std::cerr << "ERROR: --engine requires option argument" << std::endl;
return false;
}
break;
case 'b':
if (optarg)
{
args.batchSize.push_back(std::stoi(optarg));
}
else
{
std::cerr << "ERROR: --batch_size requires option argument" << std::endl;
return false;
}
break;
case 's':
if (optarg)
{
args.sequenceLength = std::stoi(optarg);
}
else
{
std::cerr << "ERROR: --sequence_length requires option argument" << std::endl;
return false;
}
break;
case 'i':
if (optarg)
{
args.iterations = std::stoi(optarg);
}
else
{
std::cerr << "ERROR: --iterations requires option argument" << std::endl;
return false;
}
break;
case 'w':
if (optarg)
{
args.warmUpRuns = std::stoi(optarg);
}
else
{
std::cerr << "ERROR: --warm_up_runs requires option argument" << std::endl;
return false;
}
break;
case 'r':
if (optarg)
{
args.randomSeed = std::stoi(optarg);
}
else
{
std::cerr << "ERROR: --random_seed requires option argument" << std::endl;
return false;
}
break;
case 'g': args.enableGraph = true; break;
default: return false;
}
}
return true;
}
// clang-format off
#define gpuErrChk(ans) \
{ \
gpuAssert((ans), __FILE__, __LINE__); \
}
// clang-format on
inline void gpuAssert(cudaError_t code, const char* file, int line, bool abort = true)
{
if (code != cudaSuccess)
{
gLogError << "GPUassert: " << cudaGetErrorString(code) << " " << file << " " << line << "\n";
if (abort)
{
exit(code);
}
}
}
template <typename T>
struct TrtDestroyer
{
void operator()(T* t)
{
t->destroy();
}
};
template <typename T>
using TrtUniquePtr = std::unique_ptr<T, TrtDestroyer<T>>;
#endif // INFER_C_COMMON_H
+61
View File
@@ -0,0 +1,61 @@
/*
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include "bert_infer.h"
#include <pybind11/numpy.h>
#include <pybind11/pybind11.h>
namespace py = pybind11;
struct BertInferenceRunner
{
BertInferenceRunner(
const std::string& enginePath, const int maxBatchSize, const int maxSeqLength, const bool enableGraph)
: bert{enginePath, maxBatchSize, maxSeqLength, enableGraph}
{
}
void prepare(const int batchSize)
{
bert.prepare(0, batchSize);
}
py::array_t<float> run(py::array_t<int> inputIds, py::array_t<int> segmentIds, py::array_t<int> inputMask)
{
const void* inputIdsPtr = inputIds.request().ptr;
const void* segmentIdsPtr = segmentIds.request().ptr;
const void* inputMaskPtr = inputMask.request().ptr;
bert.run(inputIdsPtr, segmentIdsPtr, inputMaskPtr, 0, 1);
auto output = py::array_t<float>(bert.mOutputDims, (float*) bert.mHostOutput.data());
return output;
}
BertInference bert;
};
PYBIND11_MODULE(infer_c, m)
{
m.doc() = "Pybind11 plugin for Bert inference";
py::class_<BertInferenceRunner>(m, "bert_inf")
.def(py::init<const std::string&, const int, const int, const bool>())
.def("prepare", &BertInferenceRunner::prepare)
.def("run", &BertInferenceRunner::run);
}
+19
View File
@@ -0,0 +1,19 @@
/*
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include "logging.h"
Logger gLogger(Severity::kINFO);
+85
View File
@@ -0,0 +1,85 @@
/*
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef INFER_C_LOGGING_H
#define INFER_C_LOGGING_H
#include <NvInfer.h>
#include <cassert>
#include <iostream>
#include <map>
using namespace nvinfer1;
using Severity = nvinfer1::ILogger::Severity;
class Logger : public ILogger
{
public:
Logger(Severity severity)
: mOstream(&std::cout)
, mReportableSeverity(severity)
{
}
template <typename T>
Logger& operator<<(T const& obj)
{
if (mOstream != nullptr)
{
*mOstream << obj;
}
return *this;
}
Logger& report(Severity severity, const char* msg)
{
if (severity <= mReportableSeverity)
{
const std::map<Severity, std::string> prefixMapping = {{Severity::kINTERNAL_ERROR, "[DemoBERT][F] "},
{Severity::kERROR, "[DemoBERT][E] "}, {Severity::kWARNING, "[DemoBERT][W] "},
{Severity::kINFO, "[DemoBERT][I] "}, {Severity::kVERBOSE, "[DemoBERT][V] "}};
assert(prefixMapping.find(severity) != prefixMapping.end());
mOstream = &std::cout;
*this << prefixMapping.at(severity) << msg;
return *this;
}
mOstream = nullptr;
return *this;
}
private:
void log(Severity severity, const char* msg) override
{
report(severity, msg) << "\n";
}
std::ostream* mOstream;
Severity mReportableSeverity;
};
extern Logger gLogger;
#define gLogFatal gLogger.report(Severity::kINTERNAL_ERROR, "")
#define gLogError gLogger.report(Severity::kERROR, "")
#define gLogWarning gLogger.report(Severity::kWARNING, "")
#define gLogInfo gLogger.report(Severity::kINFO, "")
#define gLogVerbose gLogger.report(Severity::kVERBOSE, "")
#endif // INFER_C_LOGGING_H
+110
View File
@@ -0,0 +1,110 @@
/*
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include "bert_infer.h"
#include "common.h"
#include <limits>
#include <random>
using namespace nvinfer1;
void printHelpInfo()
{
std::cout << "usage: ./perf [-h] [-e ENGINE] [-b BATCH_SIZE] [-s SEQUENCE_LENGTH]\n";
std::cout << " [-i ITERATIONS] [-w WARM_UP_RUNS] [-r RANDOM_SEED] [--enable_graph]\n";
std::cout << "\n";
std::cout << "BERT Inference Benchmark\n";
std::cout << "\n";
std::cout << "optional arguments:\n";
std::cout << " -h, --help show this help message and exit\n";
std::cout << " -e ENGINE, --engine ENGINE\n";
std::cout << " Path to BERT TensorRT engine\n";
std::cout << " -b BATCH_SIZE, --batch_size BATCH_SIZE\n";
std::cout << " Batch size(s) to benchmark. Can be specified multiple\n";
std::cout << " times for more than one batch size. This script\n";
std::cout << " assumes that the engine has been built with one\n";
std::cout << " optimization profile for each batch size, and that\n";
std::cout << " these profiles are in order of increasing batch size.\n";
std::cout << " -s SEQUENCE_LENGTH, --sequence_length SEQUENCE_LENGTH\n";
std::cout << " Sequence length of the BERT model\n";
std::cout << " -i ITERATIONS, --iterations ITERATIONS\n";
std::cout << " Number of iterations to run when benchmarking.\n";
std::cout << " -w WARM_UP_RUNS, --warm_up_runs WARM_UP_RUNS\n";
std::cout << " Number of iterations to run prior to benchmarking.\n";
std::cout << " -r RANDOM_SEED, --random_seed RANDOM_SEED\n";
std::cout << " Random seed.\n";
std::cout << " --enable_graph\n";
std::cout << " Enable CUDA Graph.\n";
std::cout << std::endl;
}
int main(int argc, char* argv[])
{
Args args;
const bool argsOK = parseArgs(args, argc, argv);
if (args.help)
{
printHelpInfo();
return EXIT_SUCCESS;
}
if (!argsOK)
{
std::cerr << "Invalid arguments" << std::endl;
printHelpInfo();
return EXIT_FAILURE;
}
if (args.batchSize.empty())
{
args.batchSize.push_back(1);
}
const int maxBatchSize = *std::max_element(args.batchSize.begin(), args.batchSize.end());
BertInference bert(args.engine, maxBatchSize, args.sequenceLength, args.enableGraph);
std::default_random_engine generator(args.randomSeed);
std::uniform_int_distribution<int> distribution(0, std::numeric_limits<int>::max());
const int pseudoVocabSize = 30522;
const int pseudoTypeVocabSize = 2;
const int maxInputSize = args.sequenceLength * maxBatchSize;
std::vector<int> testWordIds(maxInputSize);
std::vector<int> testSegmentIds(maxInputSize);
std::vector<int> testInputMask(maxInputSize);
std::generate(
testWordIds.begin(), testWordIds.end(), [&] { return distribution(generator) % pseudoVocabSize; });
std::generate(testSegmentIds.begin(), testSegmentIds.end(),
[&] { return distribution(generator) % pseudoTypeVocabSize; });
std::generate(testInputMask.begin(), testInputMask.end(), [&] { return 1; });
for (int i = 0; i < args.batchSize.size(); i++)
{
bert.run(i, args.batchSize[i], (void*) (testWordIds.data()), (void*) (testSegmentIds.data()),
(void*) (testInputMask.data()), args.warmUpRuns, args.iterations);
}
for (int i = 0; i < args.batchSize.size(); i++)
{
bert.reportTiming(i, args.batchSize[i]);
}
return EXIT_SUCCESS;
}
+357
View File
@@ -0,0 +1,357 @@
{
"cells": [
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"# Copyright 2019 NVIDIA Corporation. All Rights Reserved.\n",
"#\n",
"# Licensed under the Apache License, Version 2.0 (the \"License\");\n",
"# you may not use this file except in compliance with the License.\n",
"# You may obtain a copy of the License at\n",
"#\n",
"# http://www.apache.org/licenses/LICENSE-2.0\n",
"#\n",
"# Unless required by applicable law or agreed to in writing, software\n",
"# distributed under the License is distributed on an \"AS IS\" BASIS,\n",
"# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.\n",
"# See the License for the specific language governing permissions and\n",
"# limitations under the License.\n",
"# =============================================================================="
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"<img src=\"https://upload.wikimedia.org/wikipedia/en/6/6d/Nvidia_image_logo.svg\" style=\"width: 90px; float: right;\">\n",
"\n",
"# QA Inference on BERT using TensorRT"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"## 1. Overview\n",
"\n",
"Bidirectional Embedding Representations from Transformers (BERT), is a method of pre-training language representations which obtains state-of-the-art results on a wide array of Natural Language Processing (NLP) tasks. \n",
"\n",
"The original paper can be found here: https://arxiv.org/abs/1810.04805.\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 1.a Learning objectives\n",
"\n",
"This notebook demonstrates:\n",
"- Inference on Question Answering (QA) task with BERT Base/Large model\n",
"- The use fine-tuned NVIDIA BERT models\n",
"- Use of BERT model with TRT"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"## 2. Requirements\n",
"\n",
"Please refer to the ReadMe file"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"## 3. BERT Inference: Question Answering\n",
"\n",
"We can run inference on a fine-tuned BERT model for tasks like Question Answering.\n",
"\n",
"Here we use a BERT model fine-tuned on a [SQuaD 2.0 Dataset](https://rajpurkar.github.io/SQuAD-explorer/) which contains 100,000+ question-answer pairs on 500+ articles combined with over 50,000 new, unanswerable questions."
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 3.a Paragraph and Queries\n",
"\n",
"The paragraph and the questions can be customized by changing the text below. Note that when using models with small sequence lengths, you should use a shorter paragraph:"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"#### Paragraph:"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"paragraph_text = \"The Apollo program, also known as Project Apollo, was the third United States human spaceflight program carried out by the National Aeronautics and Space Administration (NASA), which accomplished landing the first humans on the Moon from 1969 to 1972. First conceived during Dwight D. Eisenhower's administration as a three-man spacecraft to follow the one-man Project Mercury which put the first Americans in space, Apollo was later dedicated to President John F. Kennedy's national goal of landing a man on the Moon and returning him safely to the Earth by the end of the 1960s, which he proposed in a May 25, 1961, address to Congress. Project Mercury was followed by the two-man Project Gemini. The first manned flight of Apollo was in 1968. Apollo ran from 1961 to 1972, and was supported by the two-man Gemini program which ran concurrently with it from 1962 to 1966. Gemini missions developed some of the space travel techniques that were necessary for the success of the Apollo missions. Apollo used Saturn family rockets as launch vehicles. Apollo/Saturn vehicles were also used for an Apollo Applications Program, which consisted of Skylab, a space station that supported three manned missions in 1973-74, and the Apollo-Soyuz Test Project, a joint Earth orbit mission with the Soviet Union in 1975.\"\n",
"\n",
"# Short paragraph version for BERT models with max sequence length of 128\n",
"short_paragraph_text = \"The Apollo program was the third United States human spaceflight program. First conceived as a three-man spacecraft to follow the one-man Project Mercury which put the first Americans in space, Apollo was dedicated to President John F. Kennedy's national goal of landing a man on the Moon. The first manned flight of Apollo was in 1968. Apollo ran from 1961 to 1972 followed by the Apollo-Soyuz Test Project a joint Earth orbit mission with the Soviet Union in 1975.\""
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"#### Question:"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"question_text = \"What project put the first Americans into space?\"\n",
"#question_text = \"What year did the first manned Apollo flight occur?\"\n",
"#question_text = \"What President is credited with the original notion of putting Americans in space?\"\n",
"#question_text = \"Who did the U.S. collaborate with on an Earth orbit mission in 1975?\""
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"In this example we ask our BERT model questions related to the following paragraph:\n",
"\n",
"**The Apollo Program**\n",
"_\"The Apollo program, also known as Project Apollo, was the third United States human spaceflight program carried out by the National Aeronautics and Space Administration (NASA), which accomplished landing the first humans on the Moon from 1969 to 1972. First conceived during Dwight D. Eisenhower's administration as a three-man spacecraft to follow the one-man Project Mercury which put the first Americans in space, Apollo was later dedicated to President John F. Kennedy's national goal of landing a man on the Moon and returning him safely to the Earth by the end of the 1960s, which he proposed in a May 25, 1961, address to Congress. Project Mercury was followed by the two-man Project Gemini. The first manned flight of Apollo was in 1968. Apollo ran from 1961 to 1972, and was supported by the two-man Gemini program which ran concurrently with it from 1962 to 1966. Gemini missions developed some of the space travel techniques that were necessary for the success of the Apollo missions. Apollo used Saturn family rockets as launch vehicles. Apollo/Saturn vehicles were also used for an Apollo Applications Program, which consisted of Skylab, a space station that supported three manned missions in 1973-74, and the Apollo-Soyuz Test Project, a joint Earth orbit mission with the Soviet Union in 1975.\"_\n",
"\n",
"The questions and relative answers expected are shown below:\n",
"\n",
" - **Q1:** \"What project put the first Americans into space?\" \n",
" - **A1:** \"Project Mercury\"\n",
" - **Q2:** \"What program was created to carry out these projects and missions?\"\n",
" - **A2:** \"The Apollo program\"\n",
" - **Q3:** \"What year did the first manned Apollo flight occur?\"\n",
" - **A3:** \"1968\"\n",
" - **Q4:** \"What President is credited with the original notion of putting Americans in space?\"\n",
" - **A4:** \"John F. Kennedy\"\n",
" - **Q5:** \"Who did the U.S. collaborate with on an Earth orbit mission in 1975?\"\n",
" - **A5:** \"Soviet Union\"\n",
" - **Q6:** \"How long did Project Apollo run?\"\n",
" - **A6:** \"1961 to 1972\"\n",
" - **Q7:** \"What program helped develop space travel techniques that Project Apollo used?\"\n",
" - **A7:** \"Gemini Mission\"\n",
" - **Q8:** \"What space station supported three manned missions in 1973-1974?\"\n",
" - **A8:** \"Skylab\""
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"## Data Preprocessing\n",
"Let's convert the paragraph and the question to BERT input with the help of the tokenizer:"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"import helpers.data_processing as dp\n",
"import helpers.tokenization as tokenization\n",
"\n",
"tokenizer = tokenization.FullTokenizer(vocab_file=\"/workspace/bert/models/fine-tuned/bert_tf_v2_large_fp16_128_v2/vocab.txt\", do_lower_case=True)\n",
"\n",
"# The maximum number of tokens for the question. Questions longer than this will be truncated to this length.\n",
"max_query_length = 64\n",
"\n",
"# When splitting up a long document into chunks, how much stride to take between chunks.\n",
"doc_stride = 128\n",
"\n",
"# The maximum total input sequence length after WordPiece tokenization. \n",
"# Sequences longer than this will be truncated, and sequences shorter \n",
"max_seq_length = 128\n",
"\n",
"# Extract tokens from the paragraph\n",
"doc_tokens = dp.convert_doc_tokens(short_paragraph_text)\n",
"\n",
"# Extract features from the paragraph and question\n",
"features = dp.convert_example_to_features(doc_tokens, question_text, tokenizer, max_seq_length, doc_stride, max_query_length)\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"## TensorRT Inference"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"import tensorrt as trt\n",
"TRT_LOGGER = trt.Logger(trt.Logger.INFO)"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"import ctypes\n",
"import os\n",
"\n",
"ctypes.CDLL(\"libnvinfer_plugin.so\", mode=ctypes.RTLD_GLOBAL)"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"import pycuda.driver as cuda\n",
"import pycuda.autoinit\n",
"import collections\n",
"import numpy as np\n",
"import time\n",
"\n",
"# Load the BERT-Large Engine\n",
"with open(\"/workspace/bert/engines/bert_large_128.engine\", \"rb\") as f, \\\n",
" trt.Runtime(TRT_LOGGER) as runtime, \\\n",
" runtime.deserialize_cuda_engine(f.read()) as engine, \\\n",
" engine.create_execution_context() as context:\n",
"\n",
" # We always use batch size 1.\n",
" input_shape = (max_seq_length, 1)\n",
" input_nbytes = trt.volume(input_shape) * trt.int32.itemsize\n",
" \n",
" # Allocate device memory for inputs.\n",
" d_inputs = [cuda.mem_alloc(input_nbytes) for binding in range(3)]\n",
" # Create a stream in which to copy inputs/outputs and run inference.\n",
" stream = cuda.Stream()\n",
"\n",
" # Specify input shapes. These must be within the min/max bounds of the active profile (0th profile in this case)\n",
" # Note that input shapes can be specified on a per-inference basis, but in this case, we only have a single shape.\n",
" for binding in range(3):\n",
" context.set_binding_shape(binding, input_shape)\n",
" assert context.all_binding_shapes_specified\n",
"\n",
" # Allocate output buffer by querying the size from the context. This may be different for different input shapes.\n",
" h_output = cuda.pagelocked_empty(tuple(context.get_binding_shape(3)), dtype=np.float32)\n",
" d_output = cuda.mem_alloc(h_output.nbytes)\n",
"\n",
" print(\"\\nRunning Inference...\")\n",
"\n",
" _NetworkOutput = collections.namedtuple( # pylint: disable=invalid-name\n",
" \"NetworkOutput\",\n",
" [\"start_logits\", \"end_logits\", \"feature_index\"])\n",
" networkOutputs = []\n",
"\n",
" eval_time_elapsed = 0\n",
" for feature_index, feature in enumerate(features):\n",
" # Copy inputs\n",
" input_ids = cuda.register_host_memory(np.ascontiguousarray(feature.input_ids.ravel()))\n",
" segment_ids = cuda.register_host_memory(np.ascontiguousarray(feature.segment_ids.ravel()))\n",
" input_mask = cuda.register_host_memory(np.ascontiguousarray(feature.input_mask.ravel()))\n",
"\n",
" eval_start_time = time.time()\n",
" cuda.memcpy_htod_async(d_inputs[0], input_ids, stream)\n",
" cuda.memcpy_htod_async(d_inputs[1], segment_ids, stream)\n",
" cuda.memcpy_htod_async(d_inputs[2], input_mask, stream)\n",
"\n",
" # Run inference\n",
" context.execute_async_v2(bindings=[int(d_inp) for d_inp in d_inputs] + [int(d_output)], stream_handle=stream.handle)\n",
" # Synchronize the stream\n",
" stream.synchronize()\n",
" eval_time_elapsed += (time.time() - eval_start_time)\n",
"\n",
" # Transfer predictions back from GPU\n",
" cuda.memcpy_dtoh_async(h_output, d_output, stream)\n",
" stream.synchronize()\n",
"\n",
" for index, batch in enumerate(h_output):\n",
" # Data Post-processing\n",
" networkOutputs.append(_NetworkOutput(\n",
" start_logits = np.array(batch.squeeze()[:, 0]),\n",
" end_logits = np.array(batch.squeeze()[:, 1]),\n",
" feature_index = feature_index\n",
" ))\n",
"\n",
" eval_time_elapsed /= len(features)\n",
" \n",
" print(\"-----------------------------\")\n",
" print(\"Running Inference at {:.3f} Sentences/Sec\".format(1.0/eval_time_elapsed))\n",
" print(\"-----------------------------\")"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"## Data Post-Processing"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"Now that we have the inference results let's extract the actual answer to our question"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
" # The total number of n-best predictions to generate in the nbest_predictions.json output file\n",
" n_best_size = 20\n",
"\n",
" # The maximum length of an answer that can be generated. This is needed \n",
" # because the start and end predictions are not conditioned on one another\n",
" max_answer_length = 30\n",
"\n",
" prediction, nbest_json, scores_diff_json = dp.get_predictions(doc_tokens, features,\n",
" networkOutputs, n_best_size, max_answer_length)\n",
" \n",
" for index, output in enumerate(networkOutputs):\n",
" print(\"Processing output\")\n",
" print(\"Answer: '{}'\".format(prediction))\n",
" print(\"with prob: {:.3f}%\".format(nbest_json[0]['probability'] * 100.0))"
]
}
],
"metadata": {
"kernelspec": {
"display_name": "Python 3",
"language": "python",
"name": "python3"
},
"language_info": {
"codemirror_mode": {
"name": "ipython",
"version": 3
},
"file_extension": ".py",
"mimetype": "text/x-python",
"name": "python",
"nbconvert_exporter": "python",
"pygments_lexer": "ipython3",
"version": "3.6.9"
}
},
"nbformat": 4,
"nbformat_minor": 2
}
+254
View File
@@ -0,0 +1,254 @@
#!/usr/bin/env python3
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""
This script uses a prebuilt TensorRT BERT QA Engine to answer a question
based on the provided passage. It additionally includes an interactive mode
where multiple questions can be asked.
"""
import time
import json
import ctypes
import argparse
import collections
import numpy as np
import tensorrt as trt
import pycuda.driver as cuda
import pycuda.autoinit
import helpers.tokenization as tokenization
import helpers.data_processing as dp
TRT_LOGGER = trt.Logger(trt.Logger.INFO)
def parse_args():
"""
Parse command line arguments
"""
parser = argparse.ArgumentParser(description=__doc__)
parser.add_argument('-e', '--engine',
help='Path to BERT TensorRT engine')
parser.add_argument("-b", "--batch-size", default=1, help="Batch size for inference.", type=int)
parser.add_argument('-p', '--passage', nargs='*',
help='Text for paragraph/passage for BERT QA',
default='')
parser.add_argument('-pf', '--passage-file',
help='File containing input passage',
default='')
parser.add_argument('-q', '--question', nargs='*',
help='Text for query/question for BERT QA',
default='')
parser.add_argument('-qf', '--question-file',
help='File containing input question',
default='')
parser.add_argument('-sq', '--squad-json',
help='SQuAD json file',
default='')
parser.add_argument('-o', '--output-prediction-file',
help='Output prediction file for SQuAD evaluation',
default='./predictions.json')
parser.add_argument('-v', '--vocab-file',
help='Path to file containing entire understandable vocab')
parser.add_argument('-s', '--sequence-length',
help='The sequence length to use. Defaults to 128',
default=128, type=int)
parser.add_argument('--max-query-length',
help='The maximum length of a query in number of tokens. Queries longer than this will be truncated',
default=64, type=int)
parser.add_argument('--max-answer-length',
help='The maximum length of an answer that can be generated',
default=30, type=int)
parser.add_argument('--n-best-size',
help='Total number of n-best predictions to generate in the nbest_predictions.json output file',
default=20, type=int)
args, _ = parser.parse_known_args()
return args
if __name__ == '__main__':
args = parse_args()
paragraph_text = None
squad_examples = None
output_prediction_file = None
if not args.passage == '':
paragraph_text = ' '.join(args.passage)
elif not args.passage_file == '':
f = open(args.passage_file, 'r')
paragraph_text = f.read()
elif not args.squad_json == '':
squad_examples = dp.read_squad_json(args.squad_json)
output_prediction_file = args.output_prediction_file
else:
paragraph_text = input("Paragraph: ")
question_text = None
if not args.question == '':
question_text = ' '.join(args.question)
elif not args.question_file == '':
f = open(args.question_file, 'r')
question_text = f.read()
tokenizer = tokenization.FullTokenizer(vocab_file=args.vocab_file, do_lower_case=True)
# When splitting up a long document into chunks, how much stride to take between chunks.
doc_stride = 128
# The maximum total input sequence length after WordPiece tokenization.
# Sequences longer than this will be truncated, and sequences shorter
max_seq_length = args.sequence_length
def question_features(tokens, question):
# Extract features from the paragraph and question
return dp.convert_example_to_features(tokens, question, tokenizer, max_seq_length, doc_stride, args.max_query_length)
# Import necessary plugins for BERT TensorRT
handle = ctypes.CDLL("libnvinfer_plugin.so", mode=ctypes.RTLD_GLOBAL)
if not handle:
raise RuntimeError("Could not load plugin library. Is `libnvinfer_plugin.so` on your LD_LIBRARY_PATH?")
# The first context created will use the 0th profile. A new context must be created
# for each additional profile needed. Here, we only use batch size 1, thus we only need the first profile.
with open(args.engine, 'rb') as f, trt.Runtime(TRT_LOGGER) as runtime, \
runtime.deserialize_cuda_engine(f.read()) as engine, engine.create_execution_context() as context:
# select engine profile
selected_profile = -1
num_binding_per_profile = engine.num_bindings // engine.num_optimization_profiles
for idx in range(engine.num_optimization_profiles):
profile_shape = engine.get_profile_shape(profile_index = idx, binding = idx * num_binding_per_profile)
if profile_shape[0][1] <= args.batch_size and profile_shape[2][1] >= args.batch_size:
selected_profile = idx
break
if selected_profile == -1:
raise RuntimeError("Could not find any profile that can run batch size {}.".format(args.batch_size))
context.active_optimization_profile = selected_profile
binding_idx_offset = selected_profile * num_binding_per_profile
# Specify input shapes. These must be within the min/max bounds of the active profile
# Note that input shapes can be specified on a per-inference basis, but in this case, we only have a single shape.
input_shape = (max_seq_length, args.batch_size)
input_nbytes = trt.volume(input_shape) * trt.int32.itemsize
for binding in range(3):
context.set_binding_shape(binding_idx_offset + binding, input_shape)
assert context.all_binding_shapes_specified
# Create a stream in which to copy inputs/outputs and run inference.
stream = cuda.Stream()
# Allocate device memory for inputs.
d_inputs = [cuda.mem_alloc(input_nbytes) for binding in range(3)]
# Allocate output buffer by querying the size from the context. This may be different for different input shapes.
h_output = cuda.pagelocked_empty(tuple(context.get_binding_shape(binding_idx_offset + 3)), dtype=np.float32)
d_output = cuda.mem_alloc(h_output.nbytes)
def inference(features, tokens):
global h_output
_NetworkOutput = collections.namedtuple( # pylint: disable=invalid-name
"NetworkOutput",
["start_logits", "end_logits", "feature_index"])
networkOutputs = []
eval_time_elapsed = 0
for feature_index, feature in enumerate(features):
# Copy inputs
input_ids_batch = np.dstack([feature.input_ids] * args.batch_size).squeeze()
segment_ids_batch = np.dstack([feature.segment_ids] * args.batch_size).squeeze()
input_mask_batch = np.dstack([feature.input_mask] * args.batch_size).squeeze()
input_ids = cuda.register_host_memory(np.ascontiguousarray(input_ids_batch.ravel()))
segment_ids = cuda.register_host_memory(np.ascontiguousarray(segment_ids_batch.ravel()))
input_mask = cuda.register_host_memory(np.ascontiguousarray(input_mask_batch.ravel()))
eval_start_time = time.time()
cuda.memcpy_htod_async(d_inputs[0], input_ids, stream)
cuda.memcpy_htod_async(d_inputs[1], segment_ids, stream)
cuda.memcpy_htod_async(d_inputs[2], input_mask, stream)
# Run inference
context.execute_async_v2(bindings=[0 for i in range(binding_idx_offset)] + [int(d_inp) for d_inp in d_inputs] + [int(d_output)], stream_handle=stream.handle)
# Synchronize the stream
stream.synchronize()
eval_time_elapsed += (time.time() - eval_start_time)
# Transfer predictions back from GPU
cuda.memcpy_dtoh_async(h_output, d_output, stream)
stream.synchronize()
# Only retrieve and post-process the first batch
batch = h_output[0]
networkOutputs.append(_NetworkOutput(
start_logits = np.array(batch.squeeze()[:, 0]),
end_logits = np.array(batch.squeeze()[:, 1]),
feature_index = feature_index
))
eval_time_elapsed /= len(features)
# Total number of n-best predictions to generate in the nbest_predictions.json output file
n_best_size = 20
# The maximum length of an answer that can be generated. This is needed
# because the start and end predictions are not conditioned on one another
max_answer_length = 30
prediction, nbest_json, scores_diff_json = dp.get_predictions(tokens, features,
networkOutputs, args.n_best_size, args.max_answer_length)
return eval_time_elapsed, prediction, nbest_json
def print_single_query(eval_time_elapsed, prediction, nbest_json):
print("------------------------")
print("Running inference in {:.3f} Sentences/Sec".format(args.batch_size/eval_time_elapsed))
print("------------------------")
print("Answer: '{}'".format(prediction))
print("With probability: {:.3f}".format(nbest_json[0]['probability'] * 100.0))
if squad_examples:
all_predictions = collections.OrderedDict()
for example in squad_examples:
features = question_features(example.doc_tokens, example.question_text)
eval_time_elapsed, prediction, nbest_json = inference(features, example.doc_tokens)
all_predictions[example.id] = prediction
with open(output_prediction_file, "w") as f:
f.write(json.dumps(all_predictions, indent=4))
print("\nOutput dump to {}".format(output_prediction_file))
else:
# Extract tokecs from the paragraph
doc_tokens = dp.convert_doc_tokens(paragraph_text)
if question_text:
print("\nPassage: {}".format(paragraph_text))
print("\nQuestion: {}".format(question_text))
features = question_features(doc_tokens, question_text)
eval_time_elapsed, prediction, nbest_json = inference(features, doc_tokens)
print_single_query(eval_time_elapsed, prediction, nbest_json)
else:
# If no question text is provided, loop until the question is 'exit'
EXIT_CMDS = ["exit", "quit"]
question_text = input("Question (to exit, type one of {:}): ".format(EXIT_CMDS))
while question_text.strip() not in EXIT_CMDS:
features = question_features(doc_tokens, question_text)
eval_time_elapsed, prediction, nbest_json = inference(features, doc_tokens)
print_single_query(eval_time_elapsed, prediction, nbest_json)
question_text = input("Question (to exit, type one of {:}): ".format(EXIT_CMDS))
+214
View File
@@ -0,0 +1,214 @@
#!/usr/bin/env python3
#
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
"""
This script uses a prebuilt TensorRT BERT QA Engine to answer a question
based on the provided passage. It additionally includes an interactive mode
where multiple questions can be asked.
"""
import os
import sys
import time
import json
import argparse
import collections
import numpy as np
import helpers.tokenization as tokenization
import helpers.data_processing as dp
sys.path.insert(0, os.path.join(os.path.dirname(os.path.abspath(__file__)), 'build'))
import infer_c
def parse_args():
"""
Parse command line arguments
"""
parser = argparse.ArgumentParser(description=__doc__)
parser.add_argument('-e', '--engine',
help='Path to BERT TensorRT engine')
parser.add_argument('-p', '--passage', nargs='*',
help='Text for paragraph/passage for BERT QA',
default='')
parser.add_argument('-pf', '--passage-file',
help='File containing input passage',
default='')
parser.add_argument('-q', '--question', nargs='*',
help='Text for query/question for BERT QA',
default='')
parser.add_argument('-qf', '--question-file',
help='File containing input question',
default='')
parser.add_argument('-sq', '--squad-json',
help='SQuAD json file',
default='')
parser.add_argument('-o', '--output-prediction-file',
help='Output prediction file for SQuAD evaluation',
default='./predictions.json')
parser.add_argument('-v', '--vocab-file',
help='Path to file containing entire understandable vocab')
parser.add_argument('-s', '--sequence-length',
help='The sequence length to use. Defaults to 128',
default=128, type=int)
parser.add_argument('--max-query-length',
help='The maximum length of a query in number of tokens. Queries longer than this will be truncated',
default=64, type=int)
parser.add_argument('--max-answer-length',
help='The maximum length of an answer that can be generated',
default=30, type=int)
parser.add_argument('--n-best-size',
help='Total number of n-best predictions to generate in the nbest_predictions.json output file',
default=20, type=int)
parser.add_argument('--enable-graph',
help='Enable CUDA Graph support',
action='store_true',
default=False)
args = parser.parse_args()
return args
if __name__ == '__main__':
args = parse_args()
paragraph_text = None
squad_examples = None
output_prediction_file = None
if not args.passage == '':
paragraph_text = ' '.join(args.passage)
elif not args.passage_file == '':
f = open(args.passage_file, 'r')
paragraph_text = f.read()
elif not args.squad_json == '':
squad_examples = dp.read_squad_json(args.squad_json)
output_prediction_file = args.output_prediction_file
else:
paragraph_text = input("Paragraph: ")
question_text = None
if not args.question == '':
question_text = ' '.join(args.question)
elif not args.question_file == '':
f = open(args.question_file, 'r')
question_text = f.read()
tokenizer = tokenization.FullTokenizer(vocab_file=args.vocab_file, do_lower_case=True)
# When splitting up a long document into chunks, how much stride to take between chunks.
doc_stride = 128
# The maximum total input sequence length after WordPiece tokenization.
# Sequences longer than this will be truncated, and sequences shorter
max_seq_length = args.sequence_length
def question_features(tokens, question):
# Extract features from the paragraph and question
return dp.convert_example_to_features(tokens, question, tokenizer, max_seq_length, doc_stride, args.max_query_length)
# The first context created will use the 0th profile. A new context must be created
# for each additional profile needed. Here, we only use batch size 1, thus we only need the first profile.
# We always use batch size 1.
# Specify input shapes as (max_seq_length, 1).
# These must be within the min/max bounds of the active profile (0th profile in this case)
# Note that input shapes can be specified on a per-inference basis, but in this case, we only have a single shape.
bert = infer_c.bert_inf(args.engine, 1, max_seq_length, args.enable_graph)
bert.prepare(1)
def inference(features, tokens):
_NetworkOutput = collections.namedtuple( # pylint: disable=invalid-name
"NetworkOutput",
["start_logits", "end_logits", "feature_index"])
networkOutputs = []
eval_time_elapsed = 0
for feature_index, feature in enumerate(features):
# Copy inputs
input_ids = np.ascontiguousarray(feature.input_ids.ravel())
segment_ids = np.ascontiguousarray(feature.segment_ids.ravel())
input_mask = np.ascontiguousarray(feature.input_mask.ravel())
eval_start_time = time.time()
# Run inference
h_output = bert.run(input_ids, segment_ids, input_mask)
eval_time_elapsed += (time.time() - eval_start_time)
for index, batch in enumerate(h_output):
# Data Post-processing
networkOutputs.append(_NetworkOutput(
start_logits = np.array(batch.squeeze()[:, 0]),
end_logits = np.array(batch.squeeze()[:, 1]),
feature_index = feature_index
))
eval_time_elapsed /= len(features)
# Total number of n-best predictions to generate in the nbest_predictions.json output file
n_best_size = 20
# The maximum length of an answer that can be generated. This is needed
# because the start and end predictions are not conditioned on one another
max_answer_length = 30
prediction, nbest_json, scores_diff_json = dp.get_predictions(tokens, features,
networkOutputs, args.n_best_size, args.max_answer_length)
return eval_time_elapsed, prediction, nbest_json
def print_single_query(eval_time_elapsed, prediction, nbest_json):
print("------------------------")
print("Running inference in {:.3f} Sentences/Sec".format(1.0/eval_time_elapsed))
print("------------------------")
print("Answer: '{}'".format(prediction))
print("With probability: {:.3f}".format(nbest_json[0]['probability'] * 100.0))
if squad_examples:
all_predictions = collections.OrderedDict()
for example_index, example in enumerate(squad_examples):
print("Processing example {} of {}".format(example_index+1, len(squad_examples)), end="\r")
features = question_features(example.doc_tokens, example.question_text)
eval_time_elapsed, prediction, nbest_json = inference(features, example.doc_tokens)
all_predictions[example.id] = prediction
with open(output_prediction_file, "w") as f:
f.write(json.dumps(all_predictions, indent=4))
print("\nOutput dump to {}".format(output_prediction_file))
else:
# Extract tokecs from the paragraph
doc_tokens = dp.convert_doc_tokens(paragraph_text)
if question_text:
print("\nPassage: {}".format(paragraph_text))
print("\nQuestion: {}".format(question_text))
features = question_features(doc_tokens, question_text)
eval_time_elapsed, prediction, nbest_json = inference(features, doc_tokens)
print_single_query(eval_time_elapsed, prediction, nbest_json)
else:
# If no question text is provided, loop until the question is 'exit'
EXIT_CMDS = ["exit", "quit"]
question_text = input("Question (to exit, type one of {:}): ".format(EXIT_CMDS))
while question_text.strip() not in EXIT_CMDS:
features = question_features(doc_tokens, question_text)
eval_time_elapsed, prediction, nbest_json = inference(features, doc_tokens)
print_single_query(eval_time_elapsed, prediction, nbest_json)
question_text = input("Question (to exit, type one of {:}): ".format(EXIT_CMDS))
+133
View File
@@ -0,0 +1,133 @@
#
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
#
import argparse
import ctypes
import numpy as np
import tensorrt as trt
import pycuda.driver as cuda
import pycuda.autoinit
import numpy as np
TRT_LOGGER = trt.Logger(trt.Logger.ERROR)
class DeviceBuffer(object):
def __init__(self, shape, dtype=trt.int32):
self.buf = cuda.mem_alloc(trt.volume(shape) * dtype.itemsize)
def binding(self):
return int(self.buf)
def free(self):
self.buf.free()
def main():
parser = argparse.ArgumentParser(description='BERT Inference Benchmark')
parser.add_argument("-e", "--engine", help='Path to BERT TensorRT engine')
parser.add_argument('-b', '--batch-size', default=[], action="append", help='Batch size(s) to benchmark. Can be specified multiple times for more than one batch size. This script assumes that the engine has been built with one optimization profile for each batch size, and that these profiles are in order of increasing batch size.', type=int)
parser.add_argument('-s', '--sequence-length', default=128, help='Sequence length of the BERT model', type=int)
parser.add_argument('-i', '--iterations', default=200, help='Number of iterations to run when benchmarking each batch size.', type=int)
parser.add_argument('-w', '--warm-up-runs', default=10, help='Number of iterations to run prior to benchmarking.', type=int)
parser.add_argument('-r', '--random-seed', required=False, default=12345, help='Random seed.', type=int)
args, _ = parser.parse_known_args()
args.batch_size = args.batch_size or [1]
# Import necessary plugins for BERT TensorRT
ctypes.CDLL("libnvinfer_plugin.so", mode=ctypes.RTLD_GLOBAL)
with open(args.engine, 'rb') as f, trt.Runtime(TRT_LOGGER) as runtime, runtime.deserialize_cuda_engine(f.read()) as engine, engine.create_execution_context() as context:
# Allocate buffers large enough to store the largest batch size
max_input_shape = (args.sequence_length, max(args.batch_size))
max_output_shape = (args.sequence_length, max(args.batch_size), 2, 1, 1)
buffers = [
DeviceBuffer(max_input_shape),
DeviceBuffer(max_input_shape),
DeviceBuffer(max_input_shape),
DeviceBuffer(max_output_shape)
]
# Prepare random input
pseudo_vocab_size = 30522
pseudo_type_vocab_size = 2
np.random.seed(args.random_seed)
test_word_ids = np.random.randint(0, pseudo_vocab_size, (args.sequence_length, max(args.batch_size)), dtype=np.int32)
test_segment_ids = np.random.randint(0, pseudo_type_vocab_size, (args.sequence_length, max(args.batch_size)), dtype=np.int32)
test_input_mask = np.ones((args.sequence_length, max(args.batch_size)), dtype=np.int32)
# Copy input h2d
cuda.memcpy_htod(buffers[0].buf, test_word_ids.ravel())
cuda.memcpy_htod(buffers[1].buf, test_segment_ids.ravel())
cuda.memcpy_htod(buffers[2].buf, test_input_mask.ravel())
num_binding_per_profile = engine.num_bindings // engine.num_optimization_profiles
bench_times = {}
for idx, batch_size in enumerate(sorted(args.batch_size)):
context.active_optimization_profile = idx
# Each profile has unique bindings
binding_idx_offset = idx * num_binding_per_profile
bindings = [0] * binding_idx_offset + [buf.binding() for buf in buffers]
shapes = {
"input_ids": (args.sequence_length, batch_size),
"segment_ids": (args.sequence_length, batch_size),
"input_mask": (args.sequence_length, batch_size),
}
for binding, shape in shapes.items():
context.set_binding_shape(engine[binding] + binding_idx_offset, shape)
assert context.all_binding_shapes_specified
# Inference
total_time = 0
start = cuda.Event()
end = cuda.Event()
stream = cuda.Stream()
# Warmup
for _ in range(args.warm_up_runs):
context.execute_async_v2(bindings=bindings, stream_handle=stream.handle)
stream.synchronize()
# Timing loop
times = []
for _ in range(args.iterations):
start.record(stream)
context.execute_async_v2(bindings=bindings, stream_handle=stream.handle)
end.record(stream)
stream.synchronize()
times.append(end.time_since(start))
# Compute average time, 95th percentile time and 99th percentile time.
bench_times[batch_size] = times
[b.free() for b in buffers]
for batch_size, times in bench_times.items():
total_time = sum(times)
avg_time = total_time / float(len(times))
times.sort()
percentile95 = times[int(len(times) * 0.95)]
percentile99 = times[int(len(times) * 0.99)]
print("Running {:} iterations with Batch Size: {:}\n\tTotal Time: {:} ms \tAverage Time: {:} ms\t95th Percentile Time: {:} ms\t99th Percentile Time: {:}".format(args.iterations, batch_size, total_time, avg_time, percentile95, percentile99))
if __name__ == '__main__':
main()
+74
View File
@@ -0,0 +1,74 @@
#!/bin/bash
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
# Setup default parameters (if no command-line parameters given)
VERSION='v2'
MODEL='large'
FT_PRECISION='fp16'
SEQ_LEN='128'
FW='tf'
while test $# -gt 0
do
case "$1" in
-h) echo "Usage: sh download_model.sh [tf|pyt] [base|large] [fp16|fp32] [128|384] [v2|v1_1]"
exit 0
;;
base) MODEL='base'
;;
large) MODEL='large'
;;
fp16) FT_PRECISION='fp16'
;;
fp32) FT_PRECISION='fp32'
;;
128) SEQ_LEN='128'
;;
384) SEQ_LEN='384'
;;
v2) VERSION='v2'
;;
v1_1) VERSION='v1_1'
;;
tf) FW='tf'
;;
pyt) FW='pyt'
;;
*) echo "Invalid argument $1...exiting"
exit 0
;;
esac
shift
done
# Prepare the download directory
mkdir -p /workspace/TensorRT/demo/BERT/models/fine-tuned
cd /workspace/TensorRT/demo/BERT/models/fine-tuned
# Download the BERT fine-tuned model
echo "Downloading BERT-${FW} ${MODEL} checkpoints with precision ${FT_PRECISION} and sequence length ${SEQ_LEN} and fine-tuned for SQuAD ${VERSION} from NGC"
if [ "${FW}" = 'tf' ]; then
ngc registry model download-version nvidia/bert_tf_${VERSION}_${MODEL}_${FT_PRECISION}_${SEQ_LEN}:2
elif [ "${FW}" = 'pyt' ]; then
if [ "${MODEL}" != 'large' ] || [ "${VERSION}" != 'v1_1' ]; then
echo "Skipping. Currently only BERT-large checkpoint fine-tuned for SQuAD v1.1 available in QAT (PyTorch) workflow."
else
ngc registry model download-version nvidia/bert_pyt_onnx_large_qa_squad11_amp_fake_quant:1
fi
else
echo "Invalid framework specified for checkpoint. Run download_model.sh -h for help."
fi
+42
View File
@@ -0,0 +1,42 @@
#!/bin/bash
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
# Setup default parameters (if no command-line parameters given)
VERSION='v1.1'
while test $# -gt 0
do
case "$1" in
-h) echo "Usage: sh download_squad.sh [v2_0|v1_1]"
exit 0
;;
v2_0) VERSION='v2.0'
;;
v1_1) VERSION='v1.1'
;;
*) echo "Invalid argument $1...exiting"
exit 0
;;
esac
shift
done
# Download the SQuAD training and dev datasets
echo "Downloading SQuAD-${VERSION} training and dev datasets"
mkdir -p /workspace/TensorRT/demo/BERT/squad
cd /workspace/TensorRT/demo/BERT/squad
wget https://rajpurkar.github.io/SQuAD-explorer/dataset/train-${VERSION}.json
wget https://rajpurkar.github.io/SQuAD-explorer/dataset/dev-${VERSION}.json
+283
View File
@@ -0,0 +1,283 @@
#!/bin/bash
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
# Usage: run_benchmark(batch_sizes, model_variant: (base/large), precision: (int8/int8-qat/fp16/fp32), sequence_length, max_batch_size, gpu_arch)
run_benchmark() {
BATCH_SIZES="${1}"
MODEL_VARIANT="${2}"
PRECISION="${3}"
SEQUENCE_LENGTH="${4}"
MAX_BATCH="${5}"
GPU_ARCH="${6}"
if [ "${PRECISION}" == "int8" ] || [ "${PRECISION}" == "int8-qat" ]; then
CHECKPOINT_PRECISION="fp16"
else
CHECKPOINT_PRECISION="${PRECISION}"
fi;
CHECKPOINTS_DIR="/workspace/TensorRT/demo/BERT/models/fine-tuned/bert_tf_v2_${MODEL_VARIANT}_${CHECKPOINT_PRECISION}_${SEQUENCE_LENGTH}_v2"
SQUAD_DIR="/workspace/TensorRT/demo/BERT/squad"
ENGINE_NAME="/workspace/TensorRT/demo/BERT/engines/bert_${MODEL_VARIANT}_${PRECISION}_bs${MAX_BATCH}_seqlen${SEQUENCE_LENGTH}_benchmark.engine"
# QAT Checkpoint - available only for BERT-Large
QAT_CHECKPOINT="/workspace/TensorRT/demo/BERT/models/fine-tuned/bert_pyt_onnx_large_qa_squad11_amp_fake_quant_v1/bert_large_v1_1_fake_quant.onnx"
CUDAGRAPH_PERFBIN="/workspace/TensorRT/demo/BERT/build/perf"
echo "==== Benchmarking BERT ${MODEL_VARIANT} ${PRECISION} SEQLEN ${SEQUENCE_LENGTH} on ${GPU_ARCH} ===="
if [ ! -f ${ENGINE_NAME} ]; then
if [ ! -d ${CHECKPOINTS_DIR} ]; then
echo "Downloading checkpoints: scripts/download_model.sh ${MODEL_VARIANT} ${CHECKPOINT_PRECISION} ${SEQUENCE_LENGTH}"
scripts/download_model.sh "${MODEL_VARIANT}" "${CHECKPOINT_PRECISION}" "${SEQUENCE_LENGTH}"
fi;
if [ "${PRECISION}" == "int8-qat" ]; then
if [ ${MODEL_VARIANT} != "large" ]; then
echo "Skipping: BERT-base not supported for int8 (QAT)"
return
fi;
if [ ! -f ${QAT_CHECKPOINT} ]; then
echo "Downloading QAT checkpoint: scripts/download_model.sh pyt v1_1 ${MODEL_VARIANT}"
scripts/download_model.sh pyt v1_1 "${MODEL_VARIANT}"
fi;
PRECISION="int8"
BUILDER_ARGS="-x ${QAT_CHECKPOINT}"
else
BUILDER_ARGS="-m ${CHECKPOINTS_DIR}/model.ckpt-8144"
fi;
BUILDER_ARGS="${BUILDER_ARGS} -o ${ENGINE_NAME} ${BATCH_SIZES} -s ${SEQUENCE_LENGTH} -c ${CHECKPOINTS_DIR} -v ${CHECKPOINTS_DIR}/vocab.txt --${PRECISION}"
if [ "${PRECISION}" == "int8" ]; then
BUILDER_ARGS="${BUILDER_ARGS} --fp16 --strict --calib-num 1"
if [ "${GPU_ARCH}" == "Ampere" ] || [ "${GPU_ARCH}" == "Turing" ]; then
BUILDER_ARGS="${BUILDER_ARGS} -iln -imh"
elif [ "${GPU_ARCH}" == "Xavier" ]; then
BUILDER_ARGS="${BUILDER_ARGS} -iln"
fi;
fi;
echo "Building engine: python3 builder.py ${BUILDER_ARGS}"
python3 builder.py ${BUILDER_ARGS}
fi;
if [ "${GPU_ARCH}" == "Ampere" ]; then
# Use more iterations for faster GPUs
NUM_ITERATIONS=2000
else
NUM_ITERATIONS=1000
fi;
if [ -f ${CUDAGRAPH_PERFBIN} ]; then
echo "Running benchmark with CUDA graph acceleration: perf ${BATCH_SIZES} -s ${SEQUENCE_LENGTH} -e ${ENGINE_NAME} -w 100 -i ${NUM_ITERATIONS} --enable_graph"
${CUDAGRAPH_PERFBIN} ${BATCH_SIZES} -s ${SEQUENCE_LENGTH} -e ${ENGINE_NAME} -w 100 -i ${NUM_ITERATIONS} --enable_graph
else
echo "Running benchmark: perf.py ${BATCH_SIZES} -s ${SEQUENCE_LENGTH} -e ${ENGINE_NAME} -w 100 -i ${NUM_ITERATIONS}"
python3 perf.py ${BATCH_SIZES} -s ${SEQUENCE_LENGTH} -e ${ENGINE_NAME} -w 100 -i ${NUM_ITERATIONS}
fi;
echo
}
arg_gpu="Volta"
arg_help=0
while [[ "$#" -gt 0 ]]; do case $1 in
--gpu) arg_gpu="$2"; shift;;
-h|--help) arg_help=1;;
*) echo "Unknown parameter passed: $1"; echo "For help type: $0 --help"; exit 1;
esac; shift; done
if [ "$arg_help" -eq "1" ]; then
echo "Usage: $0 [options]"
echo " --help or -h : Print this help menu."
echo " --gpu <arch> : GPU arch. Options: 'Volta', 'Xavier', 'Turing', 'Ampere'"
exit;
fi
mkdir -p /workspace/TensorRT/demo/BERT/engines
# BERT BASE
## INT8
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "base" "int8" "128" "32" "${arg_gpu}"
run_benchmark "-b 1" "base" "int8" "128" "1" "${arg_gpu}"
run_benchmark "-b 2" "base" "int8" "128" "2" "${arg_gpu}"
run_benchmark "-b 4" "base" "int8" "128" "4" "${arg_gpu}"
run_benchmark "-b 8" "base" "int8" "128" "8" "${arg_gpu}"
run_benchmark "-b 12" "base" "int8" "128" "12" "${arg_gpu}"
run_benchmark "-b 16" "base" "int8" "128" "16" "${arg_gpu}"
run_benchmark "-b 24" "base" "int8" "128" "24" "${arg_gpu}"
run_benchmark "-b 32" "base" "int8" "128" "32" "${arg_gpu}"
run_benchmark "-b 64" "base" "int8" "128" "64" "${arg_gpu}"
run_benchmark "-b 128" "base" "int8" "128" "128" "${arg_gpu}"
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "base" "int8" "384" "32" "${arg_gpu}"
run_benchmark "-b 1" "base" "int8" "384" "1" "${arg_gpu}"
run_benchmark "-b 2" "base" "int8" "384" "2" "${arg_gpu}"
run_benchmark "-b 4" "base" "int8" "384" "4" "${arg_gpu}"
run_benchmark "-b 8" "base" "int8" "384" "8" "${arg_gpu}"
run_benchmark "-b 12" "base" "int8" "384" "12" "${arg_gpu}"
run_benchmark "-b 16" "base" "int8" "384" "16" "${arg_gpu}"
run_benchmark "-b 24" "base" "int8" "384" "24" "${arg_gpu}"
run_benchmark "-b 32" "base" "int8" "384" "32" "${arg_gpu}"
run_benchmark "-b 64" "base" "int8" "384" "64" "${arg_gpu}"
run_benchmark "-b 128" "base" "int8" "384" "128" "${arg_gpu}"
## FP16
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "base" "fp16" "128" "32" "${arg_gpu}"
run_benchmark "-b 1" "base" "fp16" "128" "1" "${arg_gpu}"
run_benchmark "-b 2" "base" "fp16" "128" "2" "${arg_gpu}"
run_benchmark "-b 4" "base" "fp16" "128" "4" "${arg_gpu}"
run_benchmark "-b 8" "base" "fp16" "128" "8" "${arg_gpu}"
run_benchmark "-b 12" "base" "fp16" "128" "12" "${arg_gpu}"
run_benchmark "-b 16" "base" "fp16" "128" "16" "${arg_gpu}"
run_benchmark "-b 24" "base" "fp16" "128" "24" "${arg_gpu}"
run_benchmark "-b 32" "base" "fp16" "128" "32" "${arg_gpu}"
run_benchmark "-b 64" "base" "fp16" "128" "64" "${arg_gpu}"
run_benchmark "-b 128" "base" "fp16" "128" "128" "${arg_gpu}"
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "base" "fp16" "384" "32" "${arg_gpu}"
run_benchmark "-b 1" "base" "fp16" "384" "1" "${arg_gpu}"
run_benchmark "-b 2" "base" "fp16" "384" "2" "${arg_gpu}"
run_benchmark "-b 4" "base" "fp16" "384" "4" "${arg_gpu}"
run_benchmark "-b 8" "base" "fp16" "384" "8" "${arg_gpu}"
run_benchmark "-b 12" "base" "fp16" "384" "12" "${arg_gpu}"
run_benchmark "-b 16" "base" "fp16" "384" "16" "${arg_gpu}"
run_benchmark "-b 24" "base" "fp16" "384" "24" "${arg_gpu}"
run_benchmark "-b 32" "base" "fp16" "384" "32" "${arg_gpu}"
run_benchmark "-b 64" "base" "fp16" "384" "64" "${arg_gpu}"
run_benchmark "-b 128" "base" "fp16" "384" "128" "${arg_gpu}"
## FP32
#run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "base" "fp32" "128" "32" "${arg_gpu}"
run_benchmark "-b 1" "base" "fp32" "128" "1" "${arg_gpu}"
run_benchmark "-b 2" "base" "fp32" "128" "2" "${arg_gpu}"
run_benchmark "-b 4" "base" "fp32" "128" "4" "${arg_gpu}"
run_benchmark "-b 8" "base" "fp32" "128" "8" "${arg_gpu}"
run_benchmark "-b 12" "base" "fp32" "128" "12" "${arg_gpu}"
run_benchmark "-b 16" "base" "fp32" "128" "16" "${arg_gpu}"
run_benchmark "-b 24" "base" "fp32" "128" "24" "${arg_gpu}"
run_benchmark "-b 32" "base" "fp32" "128" "32" "${arg_gpu}"
run_benchmark "-b 64" "base" "fp32" "128" "64" "${arg_gpu}"
run_benchmark "-b 128" "base" "fp32" "128" "128" "${arg_gpu}"
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "base" "fp32" "384" "32" "${arg_gpu}"
run_benchmark "-b 1" "base" "fp32" "384" "1" "${arg_gpu}"
run_benchmark "-b 2" "base" "fp32" "384" "2" "${arg_gpu}"
run_benchmark "-b 4" "base" "fp32" "384" "4" "${arg_gpu}"
run_benchmark "-b 8" "base" "fp32" "384" "8" "${arg_gpu}"
run_benchmark "-b 12" "base" "fp32" "384" "12" "${arg_gpu}"
run_benchmark "-b 16" "base" "fp32" "384" "16" "${arg_gpu}"
run_benchmark "-b 24" "base" "fp32" "384" "24" "${arg_gpu}"
run_benchmark "-b 32" "base" "fp32" "384" "32" "${arg_gpu}"
run_benchmark "-b 64" "base" "fp32" "384" "64" "${arg_gpu}"
run_benchmark "-b 128" "base" "fp32" "384" "128" "${arg_gpu}"
# BERT LARGE
## INT8-QAT
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "large" "int8-qat" "128" "32" "${arg_gpu}"
run_benchmark "-b 1" "large" "int8-qat" "128" "1" "${arg_gpu}"
run_benchmark "-b 2" "large" "int8-qat" "128" "2" "${arg_gpu}"
run_benchmark "-b 4" "large" "int8-qat" "128" "4" "${arg_gpu}"
run_benchmark "-b 8" "large" "int8-qat" "128" "8" "${arg_gpu}"
run_benchmark "-b 12" "large" "int8-qat" "128" "12" "${arg_gpu}"
run_benchmark "-b 16" "large" "int8-qat" "128" "16" "${arg_gpu}"
run_benchmark "-b 24" "large" "int8-qat" "128" "24" "${arg_gpu}"
run_benchmark "-b 32" "large" "int8-qat" "128" "32" "${arg_gpu}"
run_benchmark "-b 64" "large" "int8-qat" "128" "64" "${arg_gpu}"
run_benchmark "-b 128" "large" "int8-qat" "128" "128" "${arg_gpu}"
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "large" "int8-qat" "384" "32" "${arg_gpu}"
run_benchmark "-b 1" "large" "int8-qat" "384" "1" "${arg_gpu}"
run_benchmark "-b 2" "large" "int8-qat" "384" "2" "${arg_gpu}"
run_benchmark "-b 4" "large" "int8-qat" "384" "4" "${arg_gpu}"
run_benchmark "-b 8" "large" "int8-qat" "384" "8" "${arg_gpu}"
run_benchmark "-b 12" "large" "int8-qat" "384" "12" "${arg_gpu}"
run_benchmark "-b 16" "large" "int8-qat" "384" "16" "${arg_gpu}"
run_benchmark "-b 24" "large" "int8-qat" "384" "24" "${arg_gpu}"
run_benchmark "-b 32" "large" "int8-qat" "384" "32" "${arg_gpu}"
run_benchmark "-b 64" "large" "int8-qat" "384" "64" "${arg_gpu}"
run_benchmark "-b 128" "large" "int8-qat" "384" "128" "${arg_gpu}"
## INT8
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "large" "int8" "128" "32" "${arg_gpu}"
run_benchmark "-b 1" "large" "int8" "128" "1" "${arg_gpu}"
run_benchmark "-b 2" "large" "int8" "128" "2" "${arg_gpu}"
run_benchmark "-b 4" "large" "int8" "128" "4" "${arg_gpu}"
run_benchmark "-b 8" "large" "int8" "128" "8" "${arg_gpu}"
run_benchmark "-b 12" "large" "int8" "128" "12" "${arg_gpu}"
run_benchmark "-b 16" "large" "int8" "128" "16" "${arg_gpu}"
run_benchmark "-b 24" "large" "int8" "128" "24" "${arg_gpu}"
run_benchmark "-b 32" "large" "int8" "128" "32" "${arg_gpu}"
run_benchmark "-b 64" "large" "int8" "128" "64" "${arg_gpu}"
run_benchmark "-b 128" "large" "int8" "128" "128" "${arg_gpu}"
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "large" "int8" "384" "32" "${arg_gpu}"
run_benchmark "-b 1" "large" "int8" "384" "1" "${arg_gpu}"
run_benchmark "-b 2" "large" "int8" "384" "2" "${arg_gpu}"
run_benchmark "-b 4" "large" "int8" "384" "4" "${arg_gpu}"
run_benchmark "-b 8" "large" "int8" "384" "8" "${arg_gpu}"
run_benchmark "-b 12" "large" "int8" "384" "12" "${arg_gpu}"
run_benchmark "-b 16" "large" "int8" "384" "16" "${arg_gpu}"
run_benchmark "-b 24" "large" "int8" "384" "24" "${arg_gpu}"
run_benchmark "-b 32" "large" "int8" "384" "32" "${arg_gpu}"
run_benchmark "-b 64" "large" "int8" "384" "64" "${arg_gpu}"
run_benchmark "-b 128" "large" "int8" "384" "128" "${arg_gpu}"
## FP16
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "large" "fp16" "128" "32" "${arg_gpu}"
run_benchmark "-b 1" "large" "fp16" "128" "1" "${arg_gpu}"
run_benchmark "-b 2" "large" "fp16" "128" "2" "${arg_gpu}"
run_benchmark "-b 4" "large" "fp16" "128" "4" "${arg_gpu}"
run_benchmark "-b 8" "large" "fp16" "128" "8" "${arg_gpu}"
run_benchmark "-b 12" "large" "fp16" "128" "12" "${arg_gpu}"
run_benchmark "-b 16" "large" "fp16" "128" "16" "${arg_gpu}"
run_benchmark "-b 24" "large" "fp16" "128" "24" "${arg_gpu}"
run_benchmark "-b 32" "large" "fp16" "128" "32" "${arg_gpu}"
run_benchmark "-b 64" "large" "fp16" "128" "64" "${arg_gpu}"
run_benchmark "-b 128" "large" "fp16" "128" "128" "${arg_gpu}"
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "large" "fp16" "384" "32" "${arg_gpu}"
run_benchmark "-b 1" "large" "fp16" "384" "1" "${arg_gpu}"
run_benchmark "-b 2" "large" "fp16" "384" "2" "${arg_gpu}"
run_benchmark "-b 4" "large" "fp16" "384" "4" "${arg_gpu}"
run_benchmark "-b 8" "large" "fp16" "384" "8" "${arg_gpu}"
run_benchmark "-b 12" "large" "fp16" "384" "12" "${arg_gpu}"
run_benchmark "-b 16" "large" "fp16" "384" "16" "${arg_gpu}"
run_benchmark "-b 24" "large" "fp16" "384" "24" "${arg_gpu}"
run_benchmark "-b 32" "large" "fp16" "384" "32" "${arg_gpu}"
run_benchmark "-b 64" "large" "fp16" "384" "64" "${arg_gpu}"
run_benchmark "-b 128" "large" "fp16" "384" "128" "${arg_gpu}"
## FP32
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "large" "fp32" "128" "32" "${arg_gpu}"
run_benchmark "-b 1" "large" "fp32" "128" "1" "${arg_gpu}"
run_benchmark "-b 2" "large" "fp32" "128" "2" "${arg_gpu}"
run_benchmark "-b 4" "large" "fp32" "128" "4" "${arg_gpu}"
run_benchmark "-b 8" "large" "fp32" "128" "8" "${arg_gpu}"
run_benchmark "-b 12" "large" "fp32" "128" "12" "${arg_gpu}"
run_benchmark "-b 16" "large" "fp32" "128" "16" "${arg_gpu}"
run_benchmark "-b 24" "large" "fp32" "128" "24" "${arg_gpu}"
run_benchmark "-b 32" "large" "fp32" "128" "32" "${arg_gpu}"
run_benchmark "-b 64" "large" "fp32" "128" "64" "${arg_gpu}"
run_benchmark "-b 128" "large" "fp32" "128" "128" "${arg_gpu}"
# run_benchmark "-b 1 -b 2 -b 4 -b 8 -b 12 -b 16 -b 24 -b 32" "large" "fp32" "384" "32" "${arg_gpu}"
run_benchmark "-b 1" "large" "fp32" "384" "1" "${arg_gpu}"
run_benchmark "-b 2" "large" "fp32" "384" "2" "${arg_gpu}"
run_benchmark "-b 4" "large" "fp32" "384" "4" "${arg_gpu}"
run_benchmark "-b 8" "large" "fp32" "384" "8" "${arg_gpu}"
run_benchmark "-b 12" "large" "fp32" "384" "12" "${arg_gpu}"
run_benchmark "-b 16" "large" "fp32" "384" "16" "${arg_gpu}"
run_benchmark "-b 24" "large" "fp32" "384" "24" "${arg_gpu}"
run_benchmark "-b 32" "large" "fp32" "384" "32" "${arg_gpu}"
run_benchmark "-b 64" "large" "fp32" "384" "64" "${arg_gpu}"
run_benchmark "-b 128" "large" "fp32" "384" "128" "${arg_gpu}"
+115
View File
@@ -0,0 +1,115 @@
#!/usr/bin/env python3
#
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
#
# Obtained from https://rajpurkar.github.io/SQuAD-explorer/
""" Official evaluation script for v1.1 of the SQuAD dataset. """
from __future__ import print_function
from collections import Counter
import string
import re
import argparse
import json
import sys
def normalize_answer(s):
"""Lower text and remove punctuation, articles and extra whitespace."""
def remove_articles(text):
return re.sub(r'\b(a|an|the)\b', ' ', text)
def white_space_fix(text):
return ' '.join(text.split())
def remove_punc(text):
exclude = set(string.punctuation)
return ''.join(ch for ch in text if ch not in exclude)
def lower(text):
return text.lower()
return white_space_fix(remove_articles(remove_punc(lower(s))))
def f1_score(prediction, ground_truth):
prediction_tokens = normalize_answer(prediction).split()
ground_truth_tokens = normalize_answer(ground_truth).split()
common = Counter(prediction_tokens) & Counter(ground_truth_tokens)
num_same = sum(common.values())
if num_same == 0:
return 0
precision = 1.0 * num_same / len(prediction_tokens)
recall = 1.0 * num_same / len(ground_truth_tokens)
f1 = (2 * precision * recall) / (precision + recall)
return f1
def exact_match_score(prediction, ground_truth):
return (normalize_answer(prediction) == normalize_answer(ground_truth))
def metric_max_over_ground_truths(metric_fn, prediction, ground_truths):
scores_for_ground_truths = []
for ground_truth in ground_truths:
score = metric_fn(prediction, ground_truth)
scores_for_ground_truths.append(score)
return max(scores_for_ground_truths)
def evaluate(dataset, predictions, f1_acc):
f1 = exact_match = total = 0
for article in dataset:
for paragraph in article['paragraphs']:
for qa in paragraph['qas']:
total += 1
if qa['id'] not in predictions:
message = 'Unanswered question ' + qa['id'] + \
' will receive score 0.'
print(message, file=sys.stderr)
continue
ground_truths = list(map(lambda x: x['text'], qa['answers']))
prediction = predictions[qa['id']]
exact_match += metric_max_over_ground_truths(
exact_match_score, prediction, ground_truths)
f1 += metric_max_over_ground_truths(
f1_score, prediction, ground_truths)
exact_match = 100.0 * exact_match / total
f1 = 100.0 * f1 / total
if (f1 < f1_acc - 0.5):
print("&&&& FAILED TensorRT BERT Squad Accuracy matches reference.")
else:
print("&&&& PASSED TensorRT BERT Squad Accuracy matches reference.")
return {'exact_match': exact_match, 'f1': f1}
if __name__ == '__main__':
expected_version = '1.1'
parser = argparse.ArgumentParser(
description='Evaluation for SQuAD ' + expected_version)
parser.add_argument('dataset_file', help='Dataset file')
parser.add_argument('prediction_file', help='Prediction File')
parser.add_argument('f1_acc', help='Reference Accuracy')
args = parser.parse_args()
with open(args.dataset_file) as dataset_file:
dataset_json = json.load(dataset_file)
if (dataset_json['version'] != expected_version):
print('Evaluation expects v-' + expected_version +
', but got dataset with v-' + dataset_json['version'],
file=sys.stderr)
dataset = dataset_json['data']
with open(args.prediction_file) as prediction_file:
predictions = json.load(prediction_file)
f1_acc = float(args.f1_acc)
print(json.dumps(evaluate(dataset, predictions, f1_acc)))
+54
View File
@@ -0,0 +1,54 @@
#!/usr/bin/env bash
#
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
#
arg_dockerfile=docker/ubuntu.Dockerfile
arg_imagename=tensorrt-ubuntu
arg_osversion=18.04
arg_cudaversion=11.0
arg_cudnnversion=8.0
arg_trtversion=7.1.3
arg_help=0
while [[ "$#" -gt 0 ]]; do case $1 in
--file) arg_dockerfile="$2"; shift;;
--tag) arg_imagename="$2"; shift;;
--os) arg_osversion="$2"; shift;;
--cuda) arg_cudaversion="$2"; shift;;
-h|--help) arg_help=1;;
*) echo "Unknown parameter passed: $1"; echo "For help type: $0 --help"; exit 1;
esac; shift; done
if [ "$arg_help" -eq "1" ]; then
echo "Usage: $0 [options]"
echo " --help or -h : Print this help menu."
echo " --file <dockerfile> : Docker file to use for build."
echo " --tag <imagename> : Image name for the generated container."
echo " --os <version> : OS version to use."
echo " --cuda <version> : CUDA version to use."
exit;
fi
extra_args=""
# Use RC builds for CUDA 11.0
if [ "$arg_cudaversion" = "11.0" ]; then
extra_args="$extra_args --build-arg NVCR_SUFFIX=-rc"
fi
docker_args="-f $arg_dockerfile --build-arg OS_VERSION=$arg_osversion --build-arg CUDA_VERSION=$arg_cudaversion $extra_args --build-arg uid=$(id -u) --build-arg gid=$(id -g) --tag=$arg_imagename ."
echo "Building container:"
echo "> docker build $docker_args"
docker build $docker_args
+19 -6
View File
@@ -1,4 +1,4 @@
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -12,9 +12,10 @@
# See the License for the specific language governing permissions and
# limitations under the License.
ARG CUDA_VERSION=10.2
ARG CENTOS_VERSION=7
FROM nvidia/cuda:${CUDA_VERSION}-cudnn7-devel-centos${CENTOS_VERSION}
ARG CUDA_VERSION=11.0
ARG OS_VERSION=7
ARG NVCR_SUFFIX=
FROM nvidia/cuda:${CUDA_VERSION}-devel-centos${OS_VERSION}${NVCR_SUFFIX}
LABEL maintainer="NVIDIA CORPORATION"
@@ -24,6 +25,7 @@ RUN groupadd -r -f -g ${gid} trtuser && useradd -r -u ${uid} -g ${gid} -ms /bin/
RUN usermod -aG wheel trtuser
RUN echo 'trtuser:nvidia' | chpasswd
RUN mkdir -p /workspace && chown trtuser /workspace
# Install requried libraries
RUN yum -y install \
libcurl4-openssl-dev \
@@ -34,14 +36,18 @@ RUN yum -y install \
python3 \
python3-pip \
python3-dev \
python3-setuptools \
python3-devel \
python3-wheel \
unzip \
sudo \
make
make \
build-essential
RUN cd /usr/local/bin &&\
ln -s /usr/bin/python3 python &&\
ln -s /usr/bin/pip3 pip
RUN pip3 install --upgrade pip
RUN pip3 install setuptools>=41.0.0
# Install Cmake
RUN cd /tmp && \
@@ -50,6 +56,13 @@ RUN cd /tmp && \
./cmake-3.14.4-Linux-x86_64.sh --prefix=/usr/local --exclude-subdir --skip-license && \
rm ./cmake-3.14.4-Linux-x86_64.sh
# Install PyPI packages
COPY requirements.txt /tmp/requirements.txt
RUN pip3 install -r /tmp/requirements.txt
# Download NGC client
RUN cd /usr/local/bin && wget https://ngc.nvidia.com/downloads/ngccli_bat_linux.zip && unzip ngccli_bat_linux.zip && chmod u+x ngc && rm ngccli_bat_linux.zip ngc.md5 && echo "no-apikey\nascii\nno-org\nno-team\nno-ace\n" | ngc config set
# Set environment and working directory
ENV TRT_RELEASE /tensorrt
ENV TRT_SOURCE /workspace/TensorRT
+53
View File
@@ -0,0 +1,53 @@
#!/usr/bin/env bash
#
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
#
arg_dockerfile=docker/ubuntu
arg_imagename=tensorrt-ubuntu
arg_gpus=all
arg_trtrelease=$TRT_RELEASE
arg_trtsource=$TRT_SOURCE
arg_help=0
while [[ "$#" -gt 0 ]]; do case $1 in
--tag) arg_imagename="$2"; shift;;
--gpus) arg_gpus="$2"; shift;;
--release) arg_trtrelease="$2"; shift;;
--source) arg_trtsource="$2"; shift;;
-h|--help) arg_help=1;;
*) echo "Unknown parameter passed: $1"; echo "For help type: $0 --help"; exit 1;
esac; shift; done
if [ "$arg_help" -eq "1" ]; then
echo "Usage: $0 [options]"
echo " --help or -h : Print this help menu."
echo " --tag <imagename> : Image name for the generated container."
echo " --gpus <number> : Number of GPUs visible in container. Set 'none' to disable, and 'all' to make all visible."
echo " --release <path> : Path to TensorRT release build."
echo " --source <path> : Path to TensorRT open source codebase."
exit;
fi
extra_args=""
if [ "$arg_gpus" != "none" ]; then
extra_args="$extra_args --gpus $arg_gpus"
fi
docker_args="$extra_args -v $arg_trtrelease:/tensorrt -v $arg_trtsource:/workspace/TensorRT -it $arg_imagename:latest"
echo "Launching container:"
echo "> docker run $docker_args"
docker run $docker_args
+28 -19
View File
@@ -1,4 +1,4 @@
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -12,9 +12,10 @@
# See the License for the specific language governing permissions and
# limitations under the License.
ARG CUDA_VERSION=10.0
ARG UBUNTU_VERSION=18.04
FROM nvidia/cuda:${CUDA_VERSION}-devel-ubuntu${UBUNTU_VERSION}
ARG CUDA_VERSION=10.2
ARG OS_VERSION=18.04
ARG NVCR_SUFFIX=
FROM nvidia/cuda:${CUDA_VERSION}-devel-ubuntu${OS_VERSION}${NVCR_SUFFIX}
LABEL maintainer="NVIDIA CORPORATION"
@@ -24,6 +25,7 @@ RUN groupadd -r -f -g ${gid} trtuser && useradd -r -u ${uid} -g ${gid} -ms /bin/
RUN usermod -aG sudo trtuser
RUN echo 'trtuser:nvidia' | chpasswd
RUN mkdir -p /workspace && chown trtuser /workspace
# Install requried libraries
RUN apt-get update && apt-get install -y software-properties-common
RUN add-apt-repository ppa:ubuntu-toolchain-r/test
@@ -36,18 +38,20 @@ RUN apt-get update && apt-get install -y --no-install-recommends \
python3 \
python3-pip \
python3-dev \
python3-setuptools \
python3-wheel \
sudo \
ssh \
pbzip2 \
pv \
bzip2 \
unzip
unzip \
build-essential
RUN cd /usr/local/bin &&\
ln -s /usr/bin/python3 python &&\
ln -s /usr/bin/pip3 pip
RUN pip3 install --upgrade pip
RUN pip3 install setuptools>=41.0.0
# Install Cmake
RUN cd /tmp && \
@@ -56,37 +60,42 @@ RUN cd /tmp && \
./cmake-3.14.4-Linux-x86_64.sh --prefix=/usr/local --exclude-subdir --skip-license && \
rm ./cmake-3.14.4-Linux-x86_64.sh
# Install PyPI packages
COPY requirements.txt /tmp/requirements.txt
RUN pip3 install -r /tmp/requirements.txt
# Download NGC client
RUN cd /usr/local/bin && wget https://ngc.nvidia.com/downloads/ngccli_bat_linux.zip && unzip ngccli_bat_linux.zip && chmod u+x ngc && rm ngccli_bat_linux.zip ngc.md5 && echo "no-apikey\nascii\nno-org\nno-team\nno-ace\n" | ngc config set
COPY docker/jetpack_files /pdk_files
COPY scripts/stubify.sh /pdk_files
# Install CUDA cross compile toolchain
RUN dpkg -i /pdk_files/cuda-repo-cross-aarch64-10-0-local-10.0.326_1.0-1_all.deb /pdk_files/cuda-repo-ubuntu1804-10-0-local-10.0.326-410.108_1.0-1_amd64.deb \
RUN dpkg -i /pdk_files/cuda-repo-cross-aarch64*.deb /pdk_files/cuda-repo-ubuntu*_amd64.deb \
&& apt-get update \
&& apt-get install -y cuda-cross-aarch64 \
&& rm -rf /var/lib/apt/lists/*
# Unpack cudnn
RUN dpkg -x /pdk_files/libcudnn7_7.5.0.56-1+cuda10.0_arm64.deb /pdk_files/cudnn \
&& dpkg -x /pdk_files/libcudnn7-dev_7.5.0.56-1+cuda10.0_arm64.deb /pdk_files/cudnn \
RUN dpkg -x /pdk_files/libcudnn[7-8]_*-1+cuda10.[0-9]_arm64.deb /pdk_files/cudnn \
&& dpkg -x /pdk_files/libcudnn[7-8]-dev_*-1+cuda10.[0-9]_arm64.deb /pdk_files/cudnn \
&& cd /pdk_files/cudnn/usr/include/aarch64-linux-gnu \
&& cd /pdk_files/cudnn/usr/lib/aarch64-linux-gnu \
&& ln -s libcudnn.so.7 libcudnn.so \
&& cd /pdk_files/cudnn \
&& ln -s usr/include/aarch64-linux-gnu include \
&& ln -s usr/lib/aarch64-linux-gnu lib \
&& ln -s /pdk_files/cudnn/usr/include/aarch64-linux-gnu/cudnn_v7.h /usr/include/cudnn.h
&& ln -s /pdk_files/cudnn/usr/include/aarch64-linux-gnu/cudnn_v[7-9].h /usr/include/cudnn.h
# Unpack libnvinfer
#
RUN dpkg -x /pdk_files/libnvinfer6_6.0.1-1+cuda10.0_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvinfer-dev_6.0.1-1+cuda10.0_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvparsers6_6.0.1-1+cuda10.0_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvparsers-dev_6.0.1-1+cuda10.0_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvinfer-plugin6_6.0.1-1+cuda10.0_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvinfer-plugin-dev_6.0.1-1+cuda10.0_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvonnxparsers6_6.0.1-1+cuda10.0_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvonnxparsers-dev_6.0.1-1+cuda10.0_arm64.deb /pdk_files/tensorrt
RUN dpkg -x /pdk_files/libnvinfer[0-7]_*-1+cuda10.[0-9]_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvinfer-dev_*-1+cuda10.[0-9]_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvparsers[6-7]_*-1+cuda10.[0-9]_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvparsers-dev_*-1+cuda10.[0-9]_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvinfer-plugin[6-7]_*-1+cuda10.[0-9]_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvinfer-plugin-dev_*-1+cuda10.[0-9]_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvonnxparsers[6-7]_*-1+cuda10.[0-9]_arm64.deb /pdk_files/tensorrt \
&& dpkg -x /pdk_files/libnvonnxparsers-dev_*-1+cuda10.[0-9]_arm64.deb /pdk_files/tensorrt
# create stub libraries
RUN cd /pdk_files/tensorrt \
+110
View File
@@ -0,0 +1,110 @@
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
ARG CUDA_VERSION=11.0
ARG OS_VERSION=18.04
ARG NVCR_SUFFIX=
FROM nvidia/cuda:${CUDA_VERSION}-cudnn8-devel-ubuntu${OS_VERSION}${NVCR_SUFFIX}
LABEL maintainer="NVIDIA CORPORATION"
ARG uid=1000
ARG gid=1000
RUN groupadd -r -f -g ${gid} trtuser && useradd -r -u ${uid} -g ${gid} -ms /bin/bash trtuser
RUN usermod -aG sudo trtuser
RUN echo 'trtuser:nvidia' | chpasswd
RUN mkdir -p /workspace && chown -R trtuser:trtuser /workspace
# Install requried libraries
RUN apt-get update && apt-get install -y software-properties-common
RUN add-apt-repository ppa:ubuntu-toolchain-r/test
RUN apt-get update && apt-get install -y --no-install-recommends \
libcurl4-openssl-dev \
wget \
zlib1g-dev \
git \
pkg-config \
python3 \
python3-pip \
python3-dev \
python3-setuptools \
python3-wheel \
sudo \
ssh \
pbzip2 \
pv \
bzip2 \
unzip \
g++-powerpc64le-linux-gnu \
libc6-powerpc-cross
RUN cd /usr/local/bin &&\
ln -s /usr/bin/python3 python &&\
ln -s /usr/bin/pip3 pip
RUN cd /tmp && \
wget https://github.com/Kitware/CMake/releases/download/v3.14.4/cmake-3.14.4-Linux-x86_64.sh && \
chmod +x cmake-3.14.4-Linux-x86_64.sh && \
./cmake-3.14.4-Linux-x86_64.sh --prefix=/usr/local --exclude-subdir --skip-license && \
rm ./cmake-3.14.4-Linux-x86_64.sh
# Download ppc Cudnn, Cublas, Cudart, RT, nvprof
# TODO Remove once packages are added to cuda cross compiler
RUN wget http://cuda-repo/release-candidates/Libraries/cuDNN/v8.0/8.0.2.5_20200617_28575977/11.0.x-r445/Installer/Ubuntu18_04-ppc64le/libcudnn8_8.0.2.5-1+cuda11.0_ppc64el.deb && \
wget http://cuda-repo/release-candidates/Libraries/cuDNN/v8.0/8.0.2.5_20200617_28575977/11.0.x-r445/Installer/Ubuntu18_04-ppc64le/libcudnn8-dev_8.0.2.5-1+cuda11.0_ppc64el.deb && \
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/ppc64el/libcublas-dev-11-0_11.0.0.191-1_ppc64el.deb && \
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/ppc64el/libcublas-11-0_11.0.0.191-1_ppc64el.deb && \
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/ppc64el/cuda-cudart-11-0_11.0.171-1_ppc64el.deb && \
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/ppc64el/cuda-cudart-dev-11-0_11.0.171-1_ppc64el.deb && \
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/ppc64el/cuda-nvrtc-11-0_11.0.167-1_ppc64el.deb && \
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/ppc64el/cuda-nvrtc-dev-11-0_11.0.167-1_ppc64el.deb && \
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/ppc64el/cuda-nvcc-11-0_11.0.167-1_ppc64el.deb && \
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/ppc64el/cuda-nvprof-11-0_11.0.167-1_ppc64el.deb
# Unpack Cublas
RUN dpkg -x libcublas-11-0_11.0.0.191-1_ppc64el.deb cublas && \
dpkg -x libcublas-dev-11-0_11.0.0.191-1_ppc64el.deb cublas && \
cp -r cublas/* /
# Unpack Cudart
RUN dpkg -x cuda-cudart-11-0_11.0.171-1_ppc64el.deb cudart && \
dpkg -x cuda-cudart-dev-11-0_11.0.171-1_ppc64el.deb cudart && \
cp -r cudart/* /
# Unpack RT
RUN dpkg -x cuda-nvrtc-11-0_11.0.167-1_ppc64el.deb rt && \
dpkg -x cuda-nvrtc-dev-11-0_11.0.167-1_ppc64el.deb rt && \
cp -r rt/* /
# Unpack Cudnn
RUN dpkg -x libcudnn8_8.0.2.5-1+cuda11.0_ppc64el.deb cudnn && \
dpkg -x libcudnn8-dev_8.0.2.5-1+cuda11.0_ppc64el.deb cudnn && \
cp -r cudnn/* /
# Unpack NVCC, and copy headers
RUN dpkg -x cuda-nvcc-11-0_11.0.167-1_ppc64el.deb nvcc && \
cp -r nvcc/usr/local/cuda-11.0/targets/ppc64le-linux/include/* /usr/local/cuda-11.0/targets/ppc64le-linux/include/
# Install nvprof
RUN dpkg -x cuda-nvprof-11-0_11.0.167-1_ppc64el.deb prof && \
cp -r prof/* /
# Clean up temporary files
RUN rm -rf cublas cudart rt prof nvcc
RUN rm *.deb
WORKDIR /workspace
ENV TRT_RELEASE /tensorrt
ENV TRT_SOURCE /workspace/TensorRT
USER trtuser
RUN ["/bin/bash"]
+17 -6
View File
@@ -1,4 +1,4 @@
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -12,9 +12,10 @@
# See the License for the specific language governing permissions and
# limitations under the License.
ARG CUDA_VERSION=10.2
ARG UBUNTU_VERSION=18.04
FROM nvidia/cuda:${CUDA_VERSION}-cudnn7-devel-ubuntu${UBUNTU_VERSION}
ARG CUDA_VERSION=11.0
ARG OS_VERSION=18.04
ARG NVCR_SUFFIX=
FROM nvidia/cuda:${CUDA_VERSION}-cudnn8-devel-ubuntu${OS_VERSION}${NVCR_SUFFIX}
LABEL maintainer="NVIDIA CORPORATION"
@@ -24,6 +25,7 @@ RUN groupadd -r -f -g ${gid} trtuser && useradd -r -u ${uid} -g ${gid} -ms /bin/
RUN usermod -aG sudo trtuser
RUN echo 'trtuser:nvidia' | chpasswd
RUN mkdir -p /workspace && chown trtuser /workspace
# Install requried libraries
RUN apt-get update && apt-get install -y software-properties-common
RUN add-apt-repository ppa:ubuntu-toolchain-r/test
@@ -36,18 +38,20 @@ RUN apt-get update && apt-get install -y --no-install-recommends \
python3 \
python3-pip \
python3-dev \
python3-setuptools \
python3-wheel \
sudo \
ssh \
pbzip2 \
pv \
bzip2 \
unzip
unzip \
build-essential
RUN cd /usr/local/bin &&\
ln -s /usr/bin/python3 python &&\
ln -s /usr/bin/pip3 pip
RUN pip3 install --upgrade pip
RUN pip3 install setuptools>=41.0.0
# Install Cmake
RUN cd /tmp && \
@@ -56,6 +60,13 @@ RUN cd /tmp && \
./cmake-3.14.4-Linux-x86_64.sh --prefix=/usr/local --exclude-subdir --skip-license && \
rm ./cmake-3.14.4-Linux-x86_64.sh
# Install PyPI packages
COPY requirements.txt /tmp/requirements.txt
RUN pip3 install -r /tmp/requirements.txt
# Download NGC client
RUN cd /usr/local/bin && wget https://ngc.nvidia.com/downloads/ngccli_bat_linux.zip && unzip ngccli_bat_linux.zip && chmod u+x ngc && rm ngccli_bat_linux.zip ngc.md5 && echo "no-apikey\nascii\nno-org\nno-team\nno-ace\n" | ngc config set
# Set environment and working directory
ENV TRT_RELEASE /tensorrt
ENV TRT_SOURCE /workspace/TensorRT
+7 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -19,6 +19,12 @@
#include "NvInfer.h"
//!
//! \file NvCaffeParser.h
//!
//! This is the API for the Caffe Parser
//!
//!
//! \namespace nvcaffeparser1
//!
+605 -126
View File
File diff suppressed because it is too large Load Diff
+2 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -35,6 +35,7 @@ extern "C"
//! \param libNamespace Namespace used to register all the plugins in this library
//!
TENSORRTAPI bool initLibNvInferPlugins(void* logger, const char* libNamespace);
} // extern "C"
#endif // NV_INFER_PLUGIN_H
+4 -3
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -13,12 +13,13 @@
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef NV_INFER_PLUGIN_UTILS_H
#define NV_INFER_PLUGIN_UTILS_H
#include "NvInferRuntimeCommon.h"
//!
//! \file NvPluginUtils.h
//! \file NvInferPluginUtils.h
//!
//! This is the API for the Nvidia provided TensorRT plugin utilities.
//! It lists all the parameters utilized by the TensorRT plugins.
+104 -19
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -37,14 +37,20 @@ class IPluginFactory; //!< Forward declaration of IPluginFactory for use by othe
//!
//! \brief List of supported engine capability flows.
//!
//! \note at present, kSAFE_DLA flow doesn't strictly limit execution to DLA devices - it simply
//! restricts the engine capabilities to DLA support levels anticipated in future releases.
//! The EngineCapability determines the restrictions of a network during build time for what can be executed
//! at runtime. EngineCapability::kDEFAULT does not provide any restrictions on functionality and the
//! resulting serialized engine can be executed with TensorRT's standard runtime APIs in the nvinfer1 namespace.
//! EngineCapabiltiy::kSAFE_GPU provides a restricted subset of network operations that are safety certified and
//! the resulting serialized engine can be executed with TensorRT's safe runtime APIs in the nvinfer1::safe namespace.
//! EngineCapability::kSAFE_DLA provides a restricted subset of network operations that are DLA compatible and
//! the resulting serialized engine can be executed using NvMediaDLA's runtime APIs. See sampleNvmedia for an
//! example of integrating NvMediaDLA APIs with TensorRT APIs.
//!
enum class EngineCapability : int
{
kDEFAULT = 0, //!< Full capability, TensorRT mode without any restrictions.
kSAFE_GPU = 1, //!< Safety restricted capability, TensorRT flow that can only run on GPU devices.
kSAFE_DLA = 2, //!< Safety restricted capability, TensorRT flow that can only run on DLA devices.
kDEFAULT = 0, //!< Full capability, TensorRT mode without any restrictions using TensorRT nvinfer1 APIs.
kSAFE_GPU = 1, //!< Safety restricted capability, TensorRT flow that can only run on GPU devices via TensorRT nvinfer1::safe APIs.
kSAFE_DLA = 2, //!< Safety restricted capability, TensorRT flow that can only run on DLA devices via NvMediaDLA APIs.
};
template <>
@@ -53,12 +59,15 @@ constexpr inline int EnumMax<EngineCapability>()
return 3;
} //!< Maximum number of elements in EngineCapability enum. \see EngineCapability
//!
//! \class Weights
//!
//! \brief An array of weights used as a layer parameter.
//!
//! When using the DLA, the cumulative size of all Weights used in a network
//! must be less than 512MB in size. If the build option kGPU_FALLBACK is specified,
//! then multiple DLA sub-networks may be generated from the single original network.
//!
//! The weights are held by reference until the engine has been built. Therefore the data referenced
//! by \p values field should be preserved until the build is complete.
//!
@@ -225,6 +234,8 @@ public:
//! This function is called by the implementations of INetworkDefinition, IBuilder, and ICudaEngine.
//! In particular, it is called when creating an engine and when deserializing an engine.
//!
//! \warning DataType:kBOOL not supported.
//!
virtual bool supportsFormat(DataType type, PluginFormat format) const TRTNOEXCEPT = 0;
//!
@@ -243,6 +254,8 @@ public:
//!
//! The dimensions passed here do not include the outermost batch size (i.e. for 2-D image networks, they will be 3-dimensional CHW dimensions).
//!
//! \warning DataType:kBOOL not supported.
//!
virtual void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, DataType type, PluginFormat format, int maxBatchSize) TRTNOEXCEPT = 0;
virtual ~IPluginExt() {}
@@ -677,12 +690,16 @@ public:
//! \param dlaCore The DLA core to execute the engine on (0 to N-1, where N is the maximum number of DLA's present on the device). Default value is 0.
//! \see getDLACore()
//!
//! \warning Starting with TensorRT 8, the default value will be -1 if the DLA is not specified or unused.
//!
virtual void setDLACore(int dlaCore) noexcept = 0;
//!
//! \brief Get the DLA core that the engine executes on.
//! \return If setDLACore is called, returns DLA core from 0 to N-1, else returns 0.
//!
//! \warning Starting with TensorRT 8, the default value will be -1 if the DLA is not specified or unused.
//!
virtual int getDLACore() const noexcept = 0;
//!
@@ -973,7 +990,7 @@ public:
//! then the following conditions must all hold:
//!
//! (1) minDims.nbDims == optDims.nbDims == maxDims.nbDims == networkDims.nbDims
//! (2) 1 <= minDims.d[i] <= optDims.d[i] <= maxDims.d[i] for i = 0, ..., networkDims.nbDims-1
//! (2) 0 <= minDims.d[i] <= optDims.d[i] <= maxDims.d[i] for i = 0, ..., networkDims.nbDims-1
//! (3) if networkDims.d[i] != -1, then minDims.d[i] == optDims.d[i] == maxDims.d[i] == networkDims.d[i]
//!
//! This function may (but need not be) called for an input tensor that does not have dynamic dimensions. In this
@@ -987,6 +1004,8 @@ public:
//! previously set for the same input), true if no inconsistency was detected. Note that inputs can be
//! validated only partially; a full validation is performed at engine build time.
//!
//! \warning If run on DLA, minimum, optimum, and maximum dimensions must to be the same.
//!
virtual bool setDimensions(const char* inputName, OptProfileSelector select, Dims dims) noexcept = 0;
//!
@@ -1017,6 +1036,8 @@ public:
//! tensor), else true. As for setDimensions(), a full validation can only be performed at engine build
//! time.
//!
//! \warning If run on DLA, minimum, optimum, and maximum shape values must to be the same.
//!
virtual bool setShapeValues(
const char* inputName, OptProfileSelector select, const int32_t* values, int nbValues) noexcept = 0;
@@ -1085,6 +1106,8 @@ public:
//!
//! \brief Get the number of binding indices.
//!
//! There are separate binding indices for each optimization profile.
//! This method returns the total over all profiles.
//! If the engine has been built for K profiles, the first getNbBindings() / K bindings are used by profile
//! number 0, the following getNbBindings() / K bindings are used by profile number 1 etc.
//!
@@ -1100,10 +1123,13 @@ public:
//! Engine bindings map from tensor names to indices in this array.
//! Binding indices are assigned at engine build time, and take values in the range [0 ... n-1] where n is the total number of inputs and outputs.
//!
//! To get the binding index of the name in an optimization profile with index k > 0,
//! mangle the name by appending " [profile k]", as described for method getBindingName().
//!
//! \param name The tensor name.
//! \return The binding index for the named tensor, or -1 if the name is not found.
//!
//! see getNbBindings() getBindingIndex()
//! \see getNbBindings() getBindingName()
//!
virtual int getBindingIndex(const char* name) const noexcept = 0;
@@ -1112,6 +1138,11 @@ public:
//!
//! This is the reverse mapping to that provided by getBindingIndex().
//!
//! For optimization profiles with an index k > 0, the name is mangled by appending
//! " [profile k]", with k written in decimal. For example, if the tensor in the
//! INetworkDefinition had the name "foo", and bindingIndex refers to that tensor in the
//! optimization profile with index 3, getBindingName returns "foo [profile 3]".
//!
//! \param bindingIndex The binding index.
//! \return The name corresponding to the index, or nullptr if the index is out of range.
//!
@@ -1133,8 +1164,19 @@ public:
//! \brief Get the dimensions of a binding.
//!
//! \param bindingIndex The binding index.
//! \return The dimensions of the binding if the index is in range, otherwise Dims()
//! Has -1 for any dimension with a dynamic value.
//! \return The dimensions of the binding if the index is in range, otherwise Dims().
//! Has -1 for any dimension that varies within the optimization profile.
//!
//! For example, suppose an INetworkDefinition has an input with shape [-1,-1]
//! that becomes a binding b in the engine. If the associated optimization profile
//! specifies that b has minimum dimensions as [6,9] and maximum dimensions [7,9],
//! getBindingDimensions(b) returns [-1,9], despite the second dimension being
//! dynamic in the INetworkDefinition.
//!
//! Because each optimization profile has separate bindings, the returned value can
//! differ across profiles. Consider another binding b' for the same network input,
//! but for another optimization profile. If that other profile specifies minimum
//! dimensions [5,8] and maximum dimensions [5,9], getBindingDimensions(b') returns [5,-1].
//!
//! \see getBindingIndex()
//!
@@ -1192,7 +1234,12 @@ public:
//!
//! \brief Create an execution context.
//!
//! If the engine supports dynamic shapes, each execution context in concurrent use must use a separate optimization
//! profile. The first execution context created will call setOptimizationProfile(0) implicitly. For other execution
//! contexts, setOptimizationProfile() must be called with unique profile index before calling execute or enqueue.
//!
//! \see IExecutionContext.
//! \see IExecutionContext::setOptimizationProfile()
//!
virtual IExecutionContext* createExecutionContext() noexcept = 0;
@@ -1316,13 +1363,25 @@ public:
//!
//! \brief Get the minimum / optimum / maximum dimensions for a particular binding under an optimization profile.
//!
//! \param bindingIndex The binding index (must be between 0 and getNbBindings() - 1)
//! \param bindingIndex The binding index, which must belong to the given profile,
//! or be between 0 and bindingsPerProfile-1 as described below.
//!
//! \param profileIndex The profile index (must be between 0 and getNbOptimizationProfiles()-1)
//! \param profileIndex The profile index, which must be between 0 and getNbOptimizationProfiles()-1.
//!
//! \param select Whether to query the minimum, optimum, or maximum dimensions for this binding.
//!
//! \return The minimum / optimum / maximum dimensions for this binding in this profile.
//! If the profileIndex or bindingIndex are invalid, return Dims with nbDims=-1.
//!
//! For backwards compatibility with earlier versions of TensorRT, if the bindingIndex
//! does not belong to the current optimization profile, but is between 0 and bindingsPerProfile-1,
//! where bindingsPerProfile = getNbBindings()/getNbOptimizationProfiles,
//! then a corrected bindingIndex is used instead, computed by:
//!
//! profileIndex * bindingsPerProfile + bindingIndex % bindingsPerProfile
//!
//! Otherwise the bindingIndex is considered invalid.
//!
virtual Dims getProfileDimensions(int bindingIndex, int profileIndex, OptProfileSelector select) const noexcept = 0;
//!
@@ -1340,6 +1399,12 @@ public:
//! the elementwise minimum / optimum / maximum values for this shape binding under the profile.
//! If either of the indices is out of range, or if the binding is not an input shape binding, return
//! nullptr.
//!
//! For backwards compatibility with earlier versions of TensorRT, a bindingIndex that does not belong
//! to the profile is corrected as described for getProfileDimensions.
//!
//! \see ICudaEngine::getProfileDimensions
//!
virtual const int32_t* getProfileShapeValues(int profileIndex, int inputIndex, OptProfileSelector select) const
noexcept
= 0;
@@ -1545,13 +1610,13 @@ public:
virtual const char* getName() const noexcept = 0;
//!
//! \brief set the device memory for use by this execution context.
//! \brief Set the device memory for use by this execution context.
//!
//! The memory must be aligned with cuda memory alignment property (using cudaGetDeviceProperties()), and its size must be at least that
//! returned by getDeviceMemorySize(). If using enqueue() to run the network, The memory is in
//! use from the invocation of enqueue() until network execution is complete. If using execute(),
//! it is in use until execute() returns. Releasing or otherwise using the memory for other
//! purposes during this time will result in undefined behavior.
//! The memory must be aligned with cuda memory alignment property (using cudaGetDeviceProperties()), and its size
//! must be at least that returned by getDeviceMemorySize(). Setting memory to nullptr is acceptable if
//! getDeviceMemorySize() returns 0. If using enqueue() to run the network, the memory is in use from the invocation
//! of enqueue() until network execution is complete. If using execute(), it is in use until execute() returns.
//! Releasing or otherwise using the memory for other purposes during this time will result in undefined behavior.
//!
//! \see ICudaEngine::getDeviceMemorySize() ICudaEngine::createExecutionContextWithoutDeviceMemory()
//!
@@ -1560,11 +1625,19 @@ public:
//!
//! \brief Return the strides of the buffer for the given binding.
//!
//! The strides are in units of elements, not components or bytes.
//! For example, for TensorFormat::kHWC8, a stride of one spans 8 scalars.
//!
//! Note that strides can be different for different execution contexts
//! with dynamic shapes.
//!
//! If the bindingIndex is invalid or there are dynamic dimensions that have not been
//! set yet, returns Dims with Dims::nbDims = -1.
//!
//! \param bindingIndex The binding index.
//!
//! \see getBindingComponentsPerElement
//!
virtual Dims getStrides(int bindingIndex) const noexcept = 0;
public:
@@ -1613,6 +1686,7 @@ public:
//! new dimension > 0). Furthermore, the dimensions must be in the valid range for the
//! currently selected optimization profile, and the corresponding engine must not be
//! safety-certified.
//!
//! This method will fail unless a valid optimization profile is defined for the current
//! execution context (getOptimizationProfile() must not be -1).
//!
@@ -1622,6 +1696,8 @@ public:
//!
//! \return false if an error occurs (e.g. index out of range), else true
//!
//! \see ICudaEngine::getBindingIndex
//!
virtual bool setBindingDimensions(int bindingIndex, Dims dimensions) noexcept = 0;
//!
@@ -1644,6 +1720,11 @@ public:
//!
//! \return Currently selected binding dimensions
//!
//! For backwards compatibility with earlier versions of TensorRT, a bindingIndex that does not belong
//! to the current profile is corrected as described for ICudaEngine::getProfileDimensions.
//!
//! \see ICudaEngine::getProfileDimensions
//!
virtual Dims getBindingDimensions(int bindingIndex) const noexcept = 0;
//!
@@ -1759,6 +1840,10 @@ public:
//!
//! \see ICudaEngine::getBindingIndex() ICudaEngine::getMaxBatchSize()
//!
//! \note Calling enqueueV2() with a stream in CUDA graph capture mode has a known issue. If dynamic shapes are
//! used, the first enqueueV2() call after a setInputShapeBinding() call will cause failure in stream capture
//! due to resource allocation. Please call enqueueV2() once before capturing the graph.
//!
virtual bool enqueueV2(void** bindings, cudaStream_t stream, cudaEvent_t* inputConsumed) noexcept = 0;
};
}
+45 -12
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -124,15 +124,25 @@ constexpr inline int EnumMax<ActivationType>()
//!
//! \enum DataType
//!
//! \brief The type of weights and tensors.
//!
enum class DataType : int
{
kFLOAT = 0, //!< FP32 format.
kHALF = 1, //!< FP16 format.
kINT8 = 2, //!< quantized INT8 format.
kINT32 = 3, //!< INT32 format.
kBOOL = 4 //!< BOOL format.
//! 32-bit floating point format.
kFLOAT = 0,
//! IEEE 16-bit floating-point format.
kHALF = 1,
//! 8-bit integer representing a quantized floating-point value.
kINT8 = 2,
//! Signed 32-bit integer format.
kINT32 = 3,
//! 8-bit boolean. 0 = false, 1 = true, other values undefined.
kBOOL = 4
};
template <>
@@ -170,6 +180,9 @@ constexpr inline int EnumMax<DimensionType>()
//! TensorRT can also return an invalid dims structure. This structure is represented by nbDims == -1
//! and d[i] == 0 for all d.
//!
//! TensorRT can also return an "unknown rank" dims structure. This structure is represented by nbDims == -1
//! and d[i] == -1 for all d.
//!
class Dims
{
public:
@@ -206,8 +219,11 @@ enum class TensorFormat : int
//! For a tensor with dimensions {N, C, H, W} or {numbers, channels,
//! columns, rows}, the dimensional index corresponds to {3, 2, 1, 0}
//! and thus the order is W minor.
//!
//! For DLA usage, the tensor sizes are limited to C,H,W in the range [1,8192].
//!
kLINEAR = 0,
kNCHW TRT_DEPRECATED_ENUM = kLINEAR, //! <-- Deprecated, used for backward compatibility
kNCHW TRT_DEPRECATED_ENUM = kLINEAR, //!< Deprecated name of kLINEAR, provided for backwards compatibility
//! Two wide channel vectorized row major format. This format is bound to
//! FP16. It is only available for dimensions >= 3.
@@ -216,7 +232,7 @@ enum class TensorFormat : int
//! [N][(C+1)/2][H][W][2], with the tensor coordinates (n, c, h, w)
//! mapping to array subscript [n][c/2][h][w][c%2].
kCHW2 = 1,
kNC2HW2 TRT_DEPRECATED_ENUM = kCHW2, //! <-- Deprecated, used for backward compatibility
kNC2HW2 TRT_DEPRECATED_ENUM = kCHW2, //!< Deprecated name of kCHW2, provided for backwards compatibility
//! Eight channel format where C is padded to a multiple of 8. This format
//! is bound to FP16. It is only available for dimensions >= 3.
@@ -225,14 +241,17 @@ enum class TensorFormat : int
//! [N][H][W][(C+7)/8*8], with the tensor coordinates (n, h, w, c)
//! mapping to array subscript [n][h][w][c].
kHWC8 = 2,
kNHWC8 TRT_DEPRECATED_ENUM = kHWC8, //! <-- Deprecated, used for backward compatibility
kNHWC8 TRT_DEPRECATED_ENUM = kHWC8, //!< Deprecated name of kHWC8, provided for backwards compatibility
//! Four wide channel vectorized row major format. This format is bound to
//! INT8 or FP16. It is only available for dimensions >= 3.
//! For INT8, the C dimension must be a build-time constant.
//! For a tensor with dimensions {N, C, H, W},
//! the memory layout is equivalent to a C array with dimensions
//! [N][(C+3)/4][H][W][4], with the tensor coordinates (n, c, h, w)
//! mapping to array subscript [n][c/4][h][w][c%4].
//! If running on the DLA, this format can be used for acceleration
//! with the caveat that C must equal 4.
kCHW4 = 3,
//! Sixteen wide channel vectorized row major format. This format is bound
@@ -241,6 +260,10 @@ enum class TensorFormat : int
//! the memory layout is equivalent to a C array with dimensions
//! [N][(C+15)/16][H][W][16], with the tensor coordinates (n, c, h, w)
//! mapping to array subscript [n][c/16][h][w][c%16].
//!
//! For DLA usage, this format maps to the native image format for FP16,
//! and the tensor sizes are limited to C,H,W in the range [1,8192].
//!
kCHW16 = 4,
//! Thirty-two wide channel vectorized row major format. This format is
@@ -249,6 +272,10 @@ enum class TensorFormat : int
//! the memory layout is equivalent to a C array with dimensions
//! [N][(C+31)/32][H][W][32], with the tensor coordinates (n, c, h, w)
//! mapping to array subscript [n][c/32][h][w][c%32].
//!
//! For DLA usage, this format maps to the native image format for INT8,
//! and the tensor sizes are limited to C,H,W in the range [1,8192].
//!
kCHW32 = 5
};
@@ -278,7 +305,7 @@ constexpr inline int EnumMax<TensorFormat>()
struct PluginTensorDesc
{
Dims dims;
DataType type;
DataType type; //!< \warning DataType:kBOOL not supported.
TensorFormat format;
float scale;
};
@@ -367,6 +394,8 @@ public:
//! will not be passed in, this is to keep backward compatibility with TensorRT 5.x series. Use PluginV2IOExt
//! or PluginV2DynamicExt for other PluginFormats.
//!
//! \warning DataType:kBOOL not supported.
//!
virtual bool supportsFormat(DataType type, PluginFormat format) const TRTNOEXCEPT = 0;
//!
@@ -389,6 +418,8 @@ public:
//! will not be passed in, this is to keep backward compatibility with TensorRT 5.x series. Use PluginV2IOExt
//! or PluginV2DynamicExt for other PluginFormats.
//!
//! \warning DataType:kBOOL not supported.
//!
virtual void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, DataType type, PluginFormat format, int maxBatchSize) TRTNOEXCEPT = 0;
//!
@@ -487,6 +518,8 @@ public:
//! The returned data type must have a format that is supported by the plugin.
//! \see supportsFormat()
//!
//! \warning DataType:kBOOL not supported.
//!
virtual nvinfer1::DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const TRTNOEXCEPT = 0;
//! \brief Return true if output tensor is broadcast across a batch.
@@ -724,8 +757,9 @@ enum class PluginFieldType : int
//! This information can be parsed to decode necessary plugin metadata
//!
//!
struct PluginField
class PluginField
{
public:
//!
//! \brief Plugin field attribute name
//!
@@ -1084,7 +1118,6 @@ constexpr inline int EnumMax<ErrorCode>()
return 11;
} //!< Maximum number of elements in ErrorCode enum. \see ErrorCode
//!
//! \class IErrorRecorder
//!
+12 -6
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -14,16 +14,22 @@
* limitations under the License.
*/
//!
//! \file NvInferVersion.h
//!
//! Defines the TensorRT version
//!
#ifndef NV_INFER_VERSION_H
#define NV_INFER_VERSION_H
#define NV_TENSORRT_MAJOR 7 //!< TensorRT major version.
#define NV_TENSORRT_MINOR 0 //!< TensorRT minor version.
#define NV_TENSORRT_PATCH 0 //!< TensorRT patch version.
#define NV_TENSORRT_BUILD 11 //!< TensorRT build number.
#define NV_TENSORRT_MINOR 1 //!< TensorRT minor version.
#define NV_TENSORRT_PATCH 3 //!< TensorRT patch version.
#define NV_TENSORRT_BUILD 4 //!< TensorRT build number.
#define NV_TENSORRT_SONAME_MAJOR 7 //!< Shared object library major version number.
#define NV_TENSORRT_SONAME_MINOR 0 //!< Shared object library minor version number.
#define NV_TENSORRT_SONAME_PATCH 0 //!< Shared object library patch version number.
#define NV_TENSORRT_SONAME_MINOR 1 //!< Shared object library minor version number.
#define NV_TENSORRT_SONAME_PATCH 3 //!< Shared object library patch version number.
#endif // NV_INFER_VERSION_H
+4 -4
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -14,8 +14,8 @@
* limitations under the License.
*/
#ifndef NV_ONNX_CONFIG_H
#define NV_ONNX_CONFIG_H
#ifndef NV_OnnxConfig_H
#define NV_OnnxConfig_H
#include "NvInfer.h"
@@ -185,4 +185,4 @@ TENSORRTAPI IOnnxConfig* createONNXConfig();
} // namespace nvonnxparser
#endif // NV_ONNX_CONFIG_H
#endif
+12 -6
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -21,6 +21,12 @@
#include <stddef.h>
#include <vector>
//!
//! \file NvOnnxParser.h
//!
//! This is the API for the ONNX Parser
//!
#define NV_ONNX_PARSER_MAJOR 0
#define NV_ONNX_PARSER_MINOR 1
#define NV_ONNX_PARSER_PATCH 0
@@ -116,7 +122,7 @@ public:
* This method has very limited diagnostic. If parsing the serialized model
* fails for any reason (e.g. unsupported IR version, unsupported opset, etc.)
* it the user responsibility to intercept and report the error.
* To obtain a better diagnostic, use the parseFromFile method below.
* To obtain a better diagnostic, use the parseFromFile method below.
*
* \param serialized_onnx_model Pointer to the serialized ONNX model
* \param serialized_onnx_model_size Size of the serialized ONNX model
@@ -128,14 +134,14 @@ public:
size_t serialized_onnx_model_size)
= 0;
/** \brief Parse an onnx model file, can be a binary protobuf or a text onnx model
* calls parse method inside.
/** \brief Parse an onnx model file, can be a binary protobuf or a text onnx model
* calls parse method inside.
*
* \param File name
* \param Verbosity Level
*
*
* \return true if the model was parsed successfully
*
*
*/
virtual bool parseFromFile(const char* onnxModelFile, int verbosity) = 0;
+9 -3
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -19,10 +19,16 @@
#include "NvInfer.h"
//!
//! \file NvUffParser.h
//!
//! This is the API for the UFF Parser
//!
// Current supported Universal Framework Format (UFF) version for the parser.
#define UFF_REQUIRED_VERSION_MAJOR 0
#define UFF_REQUIRED_VERSION_MINOR 6
#define UFF_REQUIRED_VERSION_PATCH 5
#define UFF_REQUIRED_VERSION_PATCH 9
//!
//! \namespace nvuffparser
@@ -277,4 +283,4 @@ TENSORRTAPI void shutdownProtobufLibrary(void) TRTNOEXCEPT;
//!
extern "C" TENSORRTAPI void* createNvUffParser_INTERNAL() TRTNOEXCEPT;
#endif /* !NV_UFF_PARSER_H */
#endif /* !NV_UFF_PARSER_H */
+7 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -19,6 +19,12 @@
#include "NvInfer.h"
//!
//! \file NvUtils.h
//!
//! This file includes various utility functions
//!
namespace nvinfer1
{
namespace utils
+3 -11
View File
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -17,23 +17,15 @@
add_custom_target(parsers DEPENDS
nvcaffeparserlibs
nvonnxparser
if (NVINTERNAL OR NVPARTNER)
nvuffparserlibs
nvparserslibs
endif()
)
add_subdirectory(caffe)
if (NVINTERNAL OR NVPARTNER)
add_subdirectory(uff)
include(uff/NvParsersCMakeLists.txt)
endif()
add_definitions("-D_PROTOBUF_INSTALL_DIR=${Protobuf_INSTALL_DIR}")
add_compile_options("-Dgoogle=google_private")
set(TENSORRT_ROOT ${PROJECT_SOURCE_DIR})
set(TENSORRT_BUILD ${TRT_BIN_DIR} ${TRT_LIB_DIR})
set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${TRT_BIN_DIR})
set(TENSORRT_BUILD ${TRT_OUT_DIR} ${TRT_LIB_DIR})
set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${TRT_OUT_DIR})
include_directories(
${Protobuf_INCLUDE_DIR}
+7 -7
View File
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -55,9 +55,9 @@ set_target_properties(${SHARED_TARGET}
CXX_STANDARD 11
CXX_STANDARD_REQUIRED YES
CXX_EXTENSIONS NO
ARCHIVE_OUTPUT_DIRECTORY "${TRT_BIN_DIR}"
LIBRARY_OUTPUT_DIRECTORY "${TRT_BIN_DIR}"
RUNTIME_OUTPUT_DIRECTORY "${TRT_BIN_DIR}"
ARCHIVE_OUTPUT_DIRECTORY "${TRT_OUT_DIR}"
LIBRARY_OUTPUT_DIRECTORY "${TRT_OUT_DIR}"
RUNTIME_OUTPUT_DIRECTORY "${TRT_OUT_DIR}"
)
target_link_libraries(${SHARED_TARGET}
@@ -105,9 +105,9 @@ set_target_properties(${STATIC_TARGET}
CXX_STANDARD 11
CXX_STANDARD_REQUIRED YES
CXX_EXTENSIONS NO
ARCHIVE_OUTPUT_DIRECTORY "${TRT_BIN_DIR}"
LIBRARY_OUTPUT_DIRECTORY "${TRT_BIN_DIR}"
RUNTIME_OUTPUT_DIRECTORY "${TRT_BIN_DIR}"
ARCHIVE_OUTPUT_DIRECTORY "${TRT_OUT_DIR}"
LIBRARY_OUTPUT_DIRECTORY "${TRT_OUT_DIR}"
RUNTIME_OUTPUT_DIRECTORY "${TRT_OUT_DIR}"
)
target_link_libraries(${STATIC_TARGET}
+1 -1
View File
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
+1 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
+1 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
+1 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
+1 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
+1 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
+1 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,4 +1,4 @@
/* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
+1 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
+1 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
+1 -1
View File
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
+9 -6
View File
@@ -51,6 +51,8 @@ set(PLUGIN_LISTS
resizeNearestPlugin
specialSlicePlugin
instanceNormalizationPlugin
groupNormalizationPlugin
coordConvACPlugin
)
# Add BERT sources if ${BERT_GENCODES} was populated
@@ -105,9 +107,9 @@ set_target_properties(${SHARED_TARGET} PROPERTIES
CXX_STANDARD "11"
CXX_STANDARD_REQUIRED "YES"
CXX_EXTENSIONS "NO"
ARCHIVE_OUTPUT_DIRECTORY "${TRT_BIN_DIR}"
LIBRARY_OUTPUT_DIRECTORY "${TRT_BIN_DIR}"
RUNTIME_OUTPUT_DIRECTORY "${TRT_BIN_DIR}"
ARCHIVE_OUTPUT_DIRECTORY "${TRT_OUT_DIR}"
LIBRARY_OUTPUT_DIRECTORY "${TRT_OUT_DIR}"
RUNTIME_OUTPUT_DIRECTORY "${TRT_OUT_DIR}"
)
set_target_properties(${SHARED_TARGET} PROPERTIES LINK_FLAGS "-Wl,--exclude-libs,ALL -Wl,--version-script=${PLUGIN_EXPORT_MAP} -Wl,--no-undefined")
@@ -124,6 +126,7 @@ target_link_libraries(${SHARED_TARGET}
${CUDART_LIB}
${CUDNN_LIB}
nvinfer
${CMAKE_DL_LIBS}
)
################################## STATIC LIBRARY #######################################
@@ -144,9 +147,9 @@ set_target_properties(${STATIC_TARGET} PROPERTIES
CXX_STANDARD "11"
CXX_STANDARD_REQUIRED "YES"
CXX_EXTENSIONS "NO"
ARCHIVE_OUTPUT_DIRECTORY "${TRT_BIN_DIR}"
LIBRARY_OUTPUT_DIRECTORY "${TRT_BIN_DIR}"
RUNTIME_OUTPUT_DIRECTORY "${TRT_BIN_DIR}"
ARCHIVE_OUTPUT_DIRECTORY "${TRT_OUT_DIR}"
LIBRARY_OUTPUT_DIRECTORY "${TRT_OUT_DIR}"
RUNTIME_OUTPUT_DIRECTORY "${TRT_OUT_DIR}"
)
set_target_properties(${STATIC_TARGET} PROPERTIES LINK_FLAGS "-Wl,--exclude-libs,ALL")
+34 -30
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
@@ -15,6 +15,7 @@
*/
#include "NvInfer.h"
#include "NvInferPlugin.h"
#include "checkMacrosPlugin.h"
#include <algorithm>
#include <array>
#include <iostream>
@@ -25,28 +26,28 @@
using namespace nvinfer1;
using namespace nvinfer1::plugin;
#include "batchedNMSPlugin/batchedNMSPlugin.h"
#include "cropAndResizePlugin/cropAndResizePlugin.h"
#include "flattenConcat/flattenConcat.h"
#include "gridAnchorPlugin/gridAnchorPlugin.h"
#include "nmsPlugin/nmsPlugin.h"
#include "normalizePlugin/normalizePlugin.h"
#include "nvFasterRCNN/nvFasterRCNNPlugin.h"
#include "priorBoxPlugin/priorBoxPlugin.h"
#include "proposalPlugin/proposalPlugin.h"
#include "regionPlugin/regionPlugin.h"
#include "reorgPlugin/reorgPlugin.h"
#include "batchTilePlugin/batchTilePlugin.h"
#include "detectionLayerPlugin/detectionLayerPlugin.h"
#include "proposalLayerPlugin/proposalLayerPlugin.h"
#include "pyramidROIAlignPlugin/pyramidROIAlignPlugin.h"
#include "resizeNearestPlugin/resizeNearestPlugin.h"
#include "specialSlicePlugin/specialSlicePlugin.h"
#include "instanceNormalizationPlugin/instanceNormalizationPlugin.h"
#include "generateDetectionPlugin/generateDetectionPlugin.h"
#include "multilevelProposeROI/multilevelProposeROIPlugin.h"
#include "multilevelCropAndResizePlugin/multilevelCropAndResizePlugin.h"
#include "batchTilePlugin.h"
#include "batchedNMSPlugin.h"
#include "coordConvACPlugin.h"
#include "cropAndResizePlugin.h"
#include "detectionLayerPlugin.h"
#include "flattenConcat.h"
#include "generateDetectionPlugin.h"
#include "gridAnchorPlugin.h"
#include "instanceNormalizationPlugin.h"
#include "multilevelCropAndResizePlugin.h"
#include "multilevelProposeROIPlugin.h"
#include "nmsPlugin.h"
#include "normalizePlugin.h"
#include "nvFasterRCNNPlugin.h"
#include "priorBoxPlugin.h"
#include "proposalLayerPlugin.h"
#include "proposalPlugin.h"
#include "pyramidROIAlignPlugin.h"
#include "regionPlugin.h"
#include "reorgPlugin.h"
#include "resizeNearestPlugin.h"
#include "specialSlicePlugin.h"
using nvinfer1::plugin::RPROIParams;
@@ -55,7 +56,8 @@ namespace nvinfer1
namespace plugin
{
ILogger* gLogger{};
extern ILogger* gLogger;
// This singleton ensures that each plugin is only registered once for a given
// namespace and type, and attempts of duplicate registration are ignored.
@@ -81,8 +83,9 @@ public:
pluginCreator->setPluginNamespace(libNamespace);
nvinfer1::plugin::gLogger = static_cast<nvinfer1::ILogger*>(logger);
std::string pluginType
= std::string(pluginCreator->getPluginNamespace()) + "::" + std::string(pluginCreator->getPluginName());
std::string pluginType = std::string{pluginCreator->getPluginNamespace()}
+ "::" + std::string{pluginCreator->getPluginName()} + " version "
+ std::string{pluginCreator->getPluginVersion()};
if (mRegistryList.find(pluginType) == mRegistryList.end())
{
@@ -91,11 +94,11 @@ public:
{
mRegistry.push(std::move(pluginCreator));
mRegistryList.insert(pluginType);
verboseMsg = "Plugin creator registration succeeded - " + pluginType;
verboseMsg = "Registered plugin creator - " + pluginType;
}
else
{
errorMsg = "Could not register plugin creator: " + pluginType;
errorMsg = "Could not register plugin creator - " + pluginType;
}
}
else
@@ -148,6 +151,7 @@ void initializePlugin(void* logger, const char* libNamespace)
} // namespace plugin
} // namespace nvinfer1
// New Plugin APIs
extern "C" {
bool initLibNvInferPlugins(void* logger, const char* libNamespace)
@@ -162,9 +166,8 @@ bool initLibNvInferPlugins(void* logger, const char* libNamespace)
initializePlugin<nvinfer1::plugin::BatchedNMSPluginCreator>(logger, libNamespace);
initializePlugin<nvinfer1::plugin::FlattenConcatPluginCreator>(logger, libNamespace);
initializePlugin<nvinfer1::plugin::CropAndResizePluginCreator>(logger, libNamespace);
initializePlugin<nvinfer1::plugin::ProposalPluginCreator>(logger, libNamespace);
initializePlugin<nvinfer1::plugin::BatchTilePluginCreator>(logger, libNamespace);
initializePlugin<nvinfer1::plugin::DetectionLayerPluginCreator>(logger, libNamespace);
initializePlugin<nvinfer1::plugin::ProposalPluginCreator>(logger, libNamespace);
initializePlugin<nvinfer1::plugin::ProposalLayerPluginCreator>(logger, libNamespace);
initializePlugin<nvinfer1::plugin::PyramidROIAlignPluginCreator>(logger, libNamespace);
initializePlugin<nvinfer1::plugin::ResizeNearestPluginCreator>(logger, libNamespace);
@@ -173,6 +176,7 @@ bool initLibNvInferPlugins(void* logger, const char* libNamespace)
initializePlugin<nvinfer1::plugin::GenerateDetectionPluginCreator>(logger, libNamespace);
initializePlugin<nvinfer1::plugin::MultilevelProposeROIPluginCreator>(logger, libNamespace);
initializePlugin<nvinfer1::plugin::MultilevelCropAndResizePluginCreator>(logger, libNamespace);
initializePlugin<nvinfer1::plugin::CoordConvACPluginCreator>(logger, libNamespace);
return true;
}
} // extern "C"
+5
View File
@@ -7,14 +7,19 @@
| [batchTilePlugin](batchTilePlugin) | BatchTilePlugin_TRT | 1 |
| [batchedNMSPlugin](batchedNMSPlugin) | BatchedNMS_TRT | 1 |
| [bertQKVToContextPlugin](bertQKVToContextPlugin) | CustomQKVToContextPluginDynamic | 1 |
| [coordConvACPlugin](coordConvACPlugin) | CoordConvAC | 1 |
| [cropAndResizePlugin](cropAndResizePlugin) | CropAndResize | 1 |
| [detectionLayerPlugin](detectionLayerPlugin) | DetectionLayer_TRT | 1 |
| [embLayerNormPlugin](embLayerNormPlugin) | CustomEmbLayerNormPluginDynamic | 1 |
| [fcPlugin](fcPlugin) | CustomFCPluginDynamic | 1 |
| [flattenConcat](flattenConcat) | FlattenConcat_TRT | 1 |
| [geluPlugin](geluPlugin) | CustomGeluPluginDynamic | 1 |
| [generateDetectionPlugin](generateDetectionPlugin) | GenerateDetection_TRT | 1 |
| [gridAnchorPlugin](gridAnchorPlugin) | GridAnchor_TRT | 1 |
| [groupNormalizationPlugin](groupNormalizationPlugin) | GroupNormalizationPlugin | 1 |
| [instanceNormalizationPlugin](instanceNormalizationPlugin) | InstanceNormalization_TRT | 1 |
| [multilevelCropAndResizePlugin](multilevelCropAndResizePlugin) | MultilevelCropAndResize_TRT | 1 |
| [multilevelProposeROI](multilevelProposeROI) | MultilevelProposeROI_TRT | 1 |
| [nmsPlugin](nmsPlugin) | NMS_TRT | 1 |
| [normalizePlugin](normalizePlugin) | Normalize_TRT | 1 |
| [nvFasterRCNN](nvFasterRCNN) | RPROI_TRT | 1 |
+1 -1
View File
@@ -1,5 +1,5 @@
#
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
+1 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
+1 -1
View File
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.

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