saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-29 20:14:01 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Merge pull request #6838 from microsoft/mzs/ortmodule-api-sync-from-master-210226

Browse source 
Sync from master
This commit is contained in:
Sherlock 2021-02-27 12:32:36 -08:00 committed by GitHub
commit 12edf22f11
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
153 changed files with 14233 additions and 12748 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

1
.gitignore vendored
View file

@ -46,6 +46,7 @@ java/gradlew.bat
java/gradle
java/.gradle
java/hs_*.log
onnxruntime/python/version_info.py
/tools/perf_util/target/classes/com/msft/send_perf_metrics
/tools/perf_util/send_perf_metrics.iml
/tools/perf_util/target/classes

1
.gitmodules vendored
View file

@ -68,4 +68,3 @@
[submodule "cmake/external/onnx-tensorrt"]
path = cmake/external/onnx-tensorrt
url = https://github.com/onnx/onnx-tensorrt.git
branch = 7.1

1283
BUILD.md
View file

File diff suppressed because it is too large Load diff

2
CONTRIBUTING.md
View file

@ -38,7 +38,7 @@ For making changes to the Windows Machine Learning WinRT API, please label your
* A feature can be implemented by you, the ONNX Runtime team, or other community members. Code contributions are greatly appreciated: feel free to work on any reviewed feature you proposed, or choose one in the backlog and send us a PR. If you are new to the project and want to work on an existing issue, we recommend starting with issues that are tagged with “good first issue”. Please let us know in the issue comments if you are actively working on implementing a feature so we can ensure it's assigned to you.
* Unit tests: New code *must* be accompanied by unit tests.
* Documentation and sample updates: If the PR affects any of the documentation or samples then include those updates in the same PR.
* Build instructions are [here](BUILD.md).
* Build instructions are [here](https://www.onnxruntime.ai/docs/how-to/build.html).
* Checkin Procedure: Once a feature is complete and tested according to the contribution guidelines follow these steps:
* Fork the repo
* git clone your fork

1
README.md
View file

@ -25,6 +25,7 @@ Common use cases for ONNX Runtime:
* [Training](https://www.onnxruntime.ai/docs/get-started/training.html)
* [Documentation](https://www.onnxruntime.ai/docs/)
* [Samples and Tutorials](https://www.onnxruntime.ai/docs/tutorials/)
* [Build Instructions](https://www.onnxruntime.ai/docs/how-to/build.html)
* [Frequently Asked Questions](./docs/FAQ.md)
## Build Pipeline Status

2
VERSION_NUMBER
View file

@ -1 +1 @@
1.6.0
1.7.0

2
cgmanifests/submodules/cgmanifest.json
View file

@ -272,7 +272,7 @@
"component": {
"type": "git",
"git": {
"commitHash": "a3a4e38b2dfa7a62b6dcae33c0d1678b3bb5ef2a",
"commitHash": "dc22bb323ece3c65419717be8a0d3d0f318a61fa",
"repositoryUrl": "https://github.com/onnx/onnx-tensorrt.git"
},
"comments": "git submodule at cmake/external/onnx-tensorrt"

25
cmake/CMakeLists.txt
View file

@ -240,7 +240,7 @@ if (onnxruntime_MINIMAL_BUILD)
if (onnxruntime_MINIMAL_BUILD_CUSTOM_OPS)
add_compile_definitions(ORT_MINIMAL_BUILD_CUSTOM_OPS)
endif()
set(onnxruntime_REDUCED_OPS_BUILD ON)
if (NOT onnxruntime_ENABLE_PYTHON)
@ -371,6 +371,12 @@ if(onnxruntime_CROSS_COMPILING)
endif()
endif()
# Mark symbols to be invisible, for macOS/iOS target only
# Due to many dependencies have different symbol visibility settings, set global compile flags here.
if (${CMAKE_SYSTEM_NAME} MATCHES "Darwin|iOS")
string(APPEND CMAKE_CXX_FLAGS " -fvisibility=hidden -fvisibility-inlines-hidden")
endif()
#must after OpenMP settings
find_package(Threads)
@ -715,7 +721,7 @@ endif()
function(onnxruntime_add_shared_library target_name)
add_library(${target_name} SHARED ${ARGN})
target_link_directories(${target_name} PRIVATE ${onnxruntime_LINK_DIRS})
target_link_directories(${target_name} PRIVATE ${onnxruntime_LINK_DIRS})
if (MSVC)
target_compile_options(${target_name} PRIVATE "$<$<COMPILE_LANGUAGE:CUDA>:SHELL:--compiler-options /utf-8>" "$<$<NOT:$<COMPILE_LANGUAGE:CUDA>>:/utf-8>")
target_compile_options(${target_name} PRIVATE "$<$<COMPILE_LANGUAGE:CUDA>:SHELL:--compiler-options /sdl>" "$<$<NOT:$<COMPILE_LANGUAGE:CUDA>>:/sdl>")
@ -753,9 +759,9 @@ function(onnxruntime_add_shared_library_module target_name)
endfunction()
#almost the same as the above function, except the first line of the body
function(onnxruntime_add_executable target_name)
function(onnxruntime_add_executable target_name)
add_executable(${target_name} ${ARGN})
target_link_directories(${target_name} PRIVATE ${onnxruntime_LINK_DIRS})
target_link_directories(${target_name} PRIVATE ${onnxruntime_LINK_DIRS})
if (MSVC)
target_compile_options(${target_name} PRIVATE "$<$<COMPILE_LANGUAGE:CUDA>:SHELL:--compiler-options /utf-8>" "$<$<NOT:$<COMPILE_LANGUAGE:CUDA>>:/utf-8>")
target_compile_options(${target_name} PRIVATE "$<$<COMPILE_LANGUAGE:CUDA>:SHELL:--compiler-options /sdl>" "$<$<NOT:$<COMPILE_LANGUAGE:CUDA>>:/sdl>")
@ -979,7 +985,7 @@ if (WIN32)
list(APPEND ORT_WARNING_FLAGS "/wd4201")
if (onnxruntime_ENABLE_STATIC_ANALYSIS)
list(APPEND ORT_WARNING_FLAGS "/analyze:stacksize 131072")
list(APPEND ORT_WARNING_FLAGS "/wd6326") # potential comparison of a constant with another constant
list(APPEND ORT_WARNING_FLAGS "/wd6326") # potential comparison of a constant with another constant
if(onnxruntime_USE_OPENMP)
list(APPEND ORT_WARNING_FLAGS "/wd6993") # Code analysis ignores OpenMP constructs
endif()
@ -1497,12 +1503,11 @@ if (onnxruntime_BUILD_NODEJS)
include(onnxruntime_nodejs.cmake)
endif()
# some of the tests rely on the shared libs to be
# built; hence the ordering
if (onnxruntime_ENABLE_PYTHON)
include(onnxruntime_python.cmake)
endif()
if (onnxruntime_BUILD_UNIT_TESTS)
if (onnxruntime_ENABLE_PYTHON)
include(onnxruntime_python.cmake)
endif()
include(onnxruntime_unittests.cmake)
endif()

2
cmake/external/dml.cmake vendored
View file

@ -20,7 +20,7 @@ if (NOT onnxruntime_USE_CUSTOM_DIRECTML)
set(NUGET_CONFIG ${PROJECT_SOURCE_DIR}/../NuGet.config)
set(PACKAGES_CONFIG ${PROJECT_SOURCE_DIR}/../packages.config)
get_filename_component(PACKAGES_DIR ${CMAKE_CURRENT_BINARY_DIR}/../packages ABSOLUTE)
set(DML_PACKAGE_DIR ${PACKAGES_DIR}/Microsoft.AI.DirectML.1.4.1)
set(DML_PACKAGE_DIR ${PACKAGES_DIR}/Microsoft.AI.DirectML.1.4.2)
set(DML_SHARED_LIB DirectML.dll)
# Restore nuget packages, which will pull down the DirectML redist package

2
cmake/external/onnx-tensorrt vendored

@ -1 +1 @@
Subproject commit b3eda616d3bb60dcd40a142e4a0a2ad95a7aa166
Subproject commit dc22bb323ece3c65419717be8a0d3d0f318a61fa

9
cmake/onnxruntime_java.cmake
View file

@ -85,6 +85,15 @@ endif()
if (onnxruntime_USE_DML)
target_compile_definitions(onnxruntime4j_jni PRIVATE USE_DIRECTML=1)
endif()
if (onnxruntime_USE_ARMNN)
target_compile_definitions(onnxruntime4j_jni PRIVATE USE_ARMNN=1)
endif()
if (onnxruntime_USE_ROCM)
target_compile_definitions(onnxruntime4j_jni PRIVATE USE_ROCM=1)
endif()
if (onnxruntime_USE_COREML)
target_compile_definitions(onnxruntime4j_jni PRIVATE USE_COREML=1)
endif()
# depend on java sources. if they change, the JNI should recompile
add_dependencies(onnxruntime4j_jni onnxruntime4j)

61
cmake/onnxruntime_providers.cmake
View file

@ -426,7 +426,7 @@ if (onnxruntime_USE_TENSORRT)
if (WIN32)
add_definitions(-D_SILENCE_EXPERIMENTAL_FILESYSTEM_DEPRECATION_WARNING=1)
set(OLD_CMAKE_CUDA_FLAGS ${CMAKE_CUDA_FLAGS})
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /wd4996 /wd4244 /wd4267 /wd4099 /wd4551 /wd4505 /wd4515 /wd4706 /wd4456 /wd4324 /wd4701 /wd4804 /wd4702")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /wd4996 /wd4244 /wd4267 /wd4099 /wd4551 /wd4505 /wd4515 /wd4706 /wd4456 /wd4324 /wd4701 /wd4804 /wd4702 /wd4458 /wd4703")
if (CMAKE_BUILD_TYPE STREQUAL "Debug")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /wd4805")
endif()
@ -568,60 +568,13 @@ if (onnxruntime_USE_OPENVINO)
)
# Header paths
list(APPEND OPENVINO_INCLUDE_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/inference_engine/include)
list(APPEND OPENVINO_INCLUDE_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/inference_engine/external/tbb/include)
list(APPEND OPENVINO_INCLUDE_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/inference_engine/external/mkltiny_lnx/include)
find_package(InferenceEngine REQUIRED)
find_package(ngraph REQUIRED)
# Library paths
list(APPEND OPENVINO_LIB_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/inference_engine/external/tbb/lib)
list(APPEND OPENVINO_LIB_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/inference_engine/external/tbb/bin)
list(APPEND OPENVINO_LIB_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/inference_engine/external/mkltiny_lnx/lib)
# Lib names
if (WIN32)
if (CMAKE_BUILD_TYPE STREQUAL "Debug")
list(APPEND OPENVINO_LIB_LIST inference_engined.lib inference_engine_legacyd.lib tbb.lib ${PYTHON_LIBRARIES})
else()
list(APPEND OPENVINO_LIB_LIST inference_engine.lib inference_engine_legacy.lib tbb.lib ${PYTHON_LIBRARIES})
endif()
if (OPENVINO_VERSION VERSION_EQUAL "2020.3")
list(APPEND OPENVINO_LIB_LIST ${InferenceEngine_LIBRARIES} ${NGRAPH_LIBRARIES} ${PYTHON_LIBRARIES})
else()
list(APPEND OPENVINO_LIB_LIST -linference_engine -linference_engine_legacy -ltbb ${PYTHON_LIBRARIES})
endif()
# Link to nGraph from OpenVINO installation
list(APPEND OPENVINO_INCLUDE_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/ngraph/include)
list(APPEND OPENVINO_INCLUDE_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/ngraph/include/ngraph/frontend)
list(APPEND OPENVINO_LIB_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/ngraph/lib)
if (OPENVINO_VERSION VERSION_EQUAL "2020.4")
if (WIN32)
list(APPEND OPENVINO_LIB_LIST ngraph.lib onnx_importer.lib)
else()
list(APPEND OPENVINO_LIB_LIST -lngraph -lonnx_importer)
endif()
endif()
if (OPENVINO_VERSION VERSION_GREATER "2020.4")
if (WIN32)
if (CMAKE_BUILD_TYPE STREQUAL "Debug")
list(APPEND OPENVINO_LIB_LIST ngraphd.lib onnx_importerd.lib)
else()
list(APPEND OPENVINO_LIB_LIST ngraph.lib onnx_importer.lib)
endif()
else()
list(APPEND OPENVINO_LIB_LIST -lngraph -lonnx_importer)
endif()
endif()
if(WIN32)
if (CMAKE_BUILD_TYPE STREQUAL "Debug")
list(APPEND OPENVINO_LIB_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/inference_engine/lib/intel64/Debug)
else()
list(APPEND OPENVINO_LIB_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/inference_engine/lib/intel64/Release)
endif()
elseif(CMAKE_SYSTEM_PROCESSOR STREQUAL "aarch64")
list(APPEND OPENVINO_LIB_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/inference_engine/lib/aarch64)
else()
list(APPEND OPENVINO_LIB_DIR_LIST $ENV{INTEL_OPENVINO_DIR}/deployment_tools/inference_engine/lib/intel64)
list(APPEND OPENVINO_LIB_LIST ${InferenceEngine_LIBRARIES} ${NGRAPH_LIBRARIES} ngraph::onnx_importer ${PYTHON_LIBRARIES})
endif()
source_group(TREE ${ONNXRUNTIME_ROOT}/core FILES ${onnxruntime_providers_openvino_cc_srcs})
@ -630,10 +583,8 @@ if (onnxruntime_USE_OPENVINO)
install(DIRECTORY ${PROJECT_SOURCE_DIR}/../include/onnxruntime/core/providers/openvino DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/onnxruntime/core/providers)
set_target_properties(onnxruntime_providers_openvino PROPERTIES LINKER_LANGUAGE CXX)
set_target_properties(onnxruntime_providers_openvino PROPERTIES FOLDER "ONNXRuntime")
target_link_directories(onnxruntime_providers_openvino PRIVATE ${OPENVINO_LIB_DIR_LIST})
add_dependencies(onnxruntime_providers_openvino onnxruntime_providers_shared ${onnxruntime_EXTERNAL_DEPENDENCIES})
target_include_directories(onnxruntime_providers_openvino SYSTEM PUBLIC ${ONNXRUNTIME_ROOT} ${CMAKE_CURRENT_BINARY_DIR} ${eigen_INCLUDE_DIRS} ${OPENVINO_INCLUDE_DIR_LIST} ${PYTHON_INCLUDE_DIRS})
# ${CMAKE_CURRENT_BINARY_DIR} is so that #include "onnxruntime_config.h" inside tensor_shape.h is found
target_link_libraries(onnxruntime_providers_openvino onnxruntime_providers_shared ${OPENVINO_LIB_LIST})
if(MSVC)

189
cmake/onnxruntime_python.cmake
View file

@ -153,6 +153,36 @@ else()
set_target_properties(onnxruntime_pybind11_state PROPERTIES SUFFIX ".so")
endif()
# Generate version_info.py in Windows build.
# Has to be done before onnxruntime_python_srcs is set.
if (WIN32)
set(VERSION_INFO_FILE "${ONNXRUNTIME_ROOT}/python/version_info.py")
if (onnxruntime_USE_CUDA)
file(WRITE "${VERSION_INFO_FILE}" "use_cuda = True\n")
file(GLOB CUDNN_DLL_PATH "${onnxruntime_CUDNN_HOME}/bin/cudnn64_*.dll")
if (NOT CUDNN_DLL_PATH)
message(FATAL_ERROR "cuDNN not found in ${onnxruntime_CUDNN_HOME}")
endif()
get_filename_component(CUDNN_DLL_NAME ${CUDNN_DLL_PATH} NAME_WE)
string(REPLACE "cudnn64_" "" CUDNN_VERSION "${CUDNN_DLL_NAME}")
file(APPEND "${VERSION_INFO_FILE}"
"cuda_version = \"${onnxruntime_CUDA_VERSION}\"\n"
"cudnn_version = \"${CUDNN_VERSION}\"\n"
)
else()
file(WRITE "${VERSION_INFO_FILE}" "use_cuda = False\n")
endif()
if ("${MSVC_TOOLSET_VERSION}" STREQUAL "142")
file(APPEND "${VERSION_INFO_FILE}" "vs2019 = True\n")
else()
file(APPEND "${VERSION_INFO_FILE}" "vs2019 = False\n")
endif()
endif()
file(GLOB onnxruntime_backend_srcs CONFIGURE_DEPENDS
"${ONNXRUNTIME_ROOT}/python/backend/*.py"
)
@ -190,20 +220,23 @@ else()
)
endif()
file(GLOB onnxruntime_python_test_srcs CONFIGURE_DEPENDS
"${ONNXRUNTIME_ROOT}/test/python/*.py"
"${ORTTRAINING_SOURCE_DIR}/test/python/*.py"
"${ORTTRAINING_SOURCE_DIR}/test/python/*.json"
)
file(GLOB onnxruntime_python_quantization_test_srcs CONFIGURE_DEPENDS
"${ONNXRUNTIME_ROOT}/test/python/quantization/*.py"
)
file(GLOB onnxruntime_python_checkpoint_test_srcs CONFIGURE_DEPENDS
"${ORTTRAINING_SOURCE_DIR}/test/python/checkpoint/*.py"
)
file(GLOB onnxruntime_python_dhp_parallel_test_srcs CONFIGURE_DEPENDS
"${ORTTRAINING_SOURCE_DIR}/test/python/dhp_parallel/*.py"
)
if (onnxruntime_BUILD_UNIT_TESTS)
file(GLOB onnxruntime_python_test_srcs CONFIGURE_DEPENDS
"${ONNXRUNTIME_ROOT}/test/python/*.py"
"${ORTTRAINING_SOURCE_DIR}/test/python/*.py"
"${ORTTRAINING_SOURCE_DIR}/test/python/*.json"
)
file(GLOB onnxruntime_python_quantization_test_srcs CONFIGURE_DEPENDS
"${ONNXRUNTIME_ROOT}/test/python/quantization/*.py"
)
file(GLOB onnxruntime_python_checkpoint_test_srcs CONFIGURE_DEPENDS
"${ORTTRAINING_SOURCE_DIR}/test/python/checkpoint/*.py"
)
file(GLOB onnxruntime_python_dhp_parallel_test_srcs CONFIGURE_DEPENDS
"${ORTTRAINING_SOURCE_DIR}/test/python/dhp_parallel/*.py"
)
endif()
file(GLOB onnxruntime_python_tools_srcs CONFIGURE_DEPENDS
"${ONNXRUNTIME_ROOT}/python/tools/*.py"
)
@ -230,109 +263,115 @@ file(GLOB onnxruntime_python_datasets_data CONFIGURE_DEPENDS
"${ONNXRUNTIME_ROOT}/python/datasets/*.onnx"
)
set(test_data_target onnxruntime_test_all)
set(build_output_target onnxruntime_common)
add_custom_command(
TARGET onnxruntime_pybind11_state POST_BUILD
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/backend
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/capi
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/capi/training
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/datasets
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/tools
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/tools/featurizer_ops
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/transformers
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/transformers/longformer
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/quantization
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/quantization/operators
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/checkpoint
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/dhp_parallel
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/quantization
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/backend
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi/training
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/datasets
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/tools
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/tools/featurizer_ops
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/transformers
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/transformers/longformer
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/quantization
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/quantization/operators
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/checkpoint
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/dhp_parallel
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/quantization
COMMAND ${CMAKE_COMMAND} -E copy
${ONNXRUNTIME_ROOT}/__init__.py
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/
COMMAND ${CMAKE_COMMAND} -E copy
${REPO_ROOT}/ThirdPartyNotices.txt
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/
COMMAND ${CMAKE_COMMAND} -E copy
${REPO_ROOT}/docs/Privacy.md
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/
COMMAND ${CMAKE_COMMAND} -E copy
${REPO_ROOT}/LICENSE
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_test_srcs}
$<TARGET_FILE_DIR:${test_data_target}>
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_quantization_test_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/quantization/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_checkpoint_test_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/checkpoint/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_dhp_parallel_test_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/dhp_parallel/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_backend_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/backend/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/backend/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/capi/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_capi_training_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/capi/training/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi/training/
COMMAND ${CMAKE_COMMAND} -E copy
$<TARGET_FILE:onnxruntime_pybind11_state>
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/capi/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_datasets_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/datasets/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/datasets/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_datasets_data}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/datasets/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/datasets/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_tools_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/tools/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/tools/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_tools_featurizers_src}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/tools/featurizer_ops/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/tools/featurizer_ops/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_quantization_src}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/quantization/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/quantization/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_quantization_operators_src}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/quantization/operators/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/quantization/operators/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_transformers_src}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/transformers/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/transformers/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_transformers_longformer_src}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/transformers/longformer/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/transformers/longformer/
COMMAND ${CMAKE_COMMAND} -E copy
${REPO_ROOT}/VERSION_NUMBER
$<TARGET_FILE_DIR:${test_data_target}>
$<TARGET_FILE_DIR:${build_output_target}>
)
if (onnxruntime_BUILD_UNIT_TESTS)
add_custom_command(
TARGET onnxruntime_pybind11_state POST_BUILD
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_test_srcs}
$<TARGET_FILE_DIR:${build_output_target}>
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_quantization_test_srcs}
$<TARGET_FILE_DIR:${build_output_target}>/quantization/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_checkpoint_test_srcs}
$<TARGET_FILE_DIR:${build_output_target}>/checkpoint/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_dhp_parallel_test_srcs}
$<TARGET_FILE_DIR:${build_output_target}>/dhp_parallel/
)
endif()
if (onnxruntime_ENABLE_TRAINING)
add_custom_command(
TARGET onnxruntime_pybind11_state POST_BUILD
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/training
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/training/amp
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/training/optim
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/training
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/training/amp
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/training/optim
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_capi_training_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/capi/training/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi/training/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_root_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/training/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/training/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_amp_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/training/amp/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/training/amp/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_optim_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/training/optim/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/training/optim/
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_train_tools_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/training/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/training/
)
endif()
@ -342,7 +381,7 @@ if (onnxruntime_USE_DNNL)
COMMAND ${CMAKE_COMMAND} -E copy
${DNNL_DLL_PATH} $<TARGET_FILE:onnxruntime_providers_dnnl>
$<TARGET_FILE:onnxruntime_providers_shared>
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/capi/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi/
)
endif()
@ -352,20 +391,18 @@ if (onnxruntime_USE_TENSORRT)
COMMAND ${CMAKE_COMMAND} -E copy
$<TARGET_FILE:onnxruntime_providers_tensorrt>
$<TARGET_FILE:onnxruntime_providers_shared>
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/capi/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi/
)
endif()
if (onnxruntime_USE_OPENVINO)
if(NOT WIN32)
add_custom_command(
TARGET onnxruntime_pybind11_state POST_BUILD
COMMAND ${CMAKE_COMMAND} -E copy
${OPENVINO_DLL_PATH} $<TARGET_FILE:onnxruntime_providers_openvino>
$<TARGET_FILE:onnxruntime_providers_openvino>
$<TARGET_FILE:onnxruntime_providers_shared>
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/capi/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi/
)
endif()
endif()
if (onnxruntime_USE_TVM)
@ -373,7 +410,7 @@ if (onnxruntime_USE_TVM)
TARGET onnxruntime_pybind11_state POST_BUILD
COMMAND ${CMAKE_COMMAND} -E copy
$<TARGET_FILE:tvm> $<TARGET_FILE:nnvm_compiler>
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/capi/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi/
)
endif()
@ -383,10 +420,10 @@ if (onnxruntime_USE_NUPHAR)
)
add_custom_command(
TARGET onnxruntime_pybind11_state POST_BUILD
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/nuphar
COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/nuphar
COMMAND ${CMAKE_COMMAND} -E copy
${onnxruntime_python_nuphar_python_srcs}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/nuphar/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/nuphar/
)
endif()
@ -395,7 +432,7 @@ if (onnxruntime_USE_DML)
TARGET onnxruntime_pybind11_state POST_BUILD
COMMAND ${CMAKE_COMMAND} -E copy
${DML_PACKAGE_DIR}/bin/${onnxruntime_target_platform}-win/${DML_SHARED_LIB}
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/capi/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi/
)
endif()
@ -404,7 +441,7 @@ if (onnxruntime_USE_NNAPI_BUILTIN)
TARGET onnxruntime_pybind11_state POST_BUILD
COMMAND ${CMAKE_COMMAND} -E copy
$<TARGET_FILE:onnxruntime_providers_nnapi>
$<TARGET_FILE_DIR:${test_data_target}>/onnxruntime/capi/
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi/
)
endif()

4
dockerfiles/Dockerfile.tensorrt
View file

@ -5,7 +5,7 @@
# Dockerfile to run ONNXRuntime with TensorRT integration
# nVidia TensorRT Base Image
FROM nvcr.io/nvidia/tensorrt:20.07.1-py3
FROM nvcr.io/nvidia/tensorrt:20.12-py3
MAINTAINER Vinitra Swamy "viswamy@microsoft.com"
ARG ONNXRUNTIME_REPO=https://github.com/Microsoft/onnxruntime
@ -26,7 +26,7 @@ RUN git clone --single-branch --branch ${ONNXRUNTIME_BRANCH} --recursive ${ONNXR
cp onnxruntime/dockerfiles/LICENSE-IMAGE.txt /code/LICENSE-IMAGE.txt &&\
cp onnxruntime/ThirdPartyNotices.txt /code/ThirdPartyNotices.txt &&\
cd onnxruntime &&\
/bin/sh ./build.sh --cuda_home /usr/local/cuda --cudnn_home /usr/lib/x86_64-linux-gnu/ --use_tensorrt --tensorrt_home /workspace/tensorrt --config Release --build_wheel --update --build --cmake_extra_defines ONNXRUNTIME_VERSION=$(cat ./VERSION_NUMBER) &&\
/bin/sh ./build.sh --parallel --cuda_home /usr/local/cuda --cudnn_home /usr/lib/x86_64-linux-gnu/ --use_tensorrt --tensorrt_home /workspace/tensorrt --config Release --build_wheel --update --build --cmake_extra_defines ONNXRUNTIME_VERSION=$(cat ./VERSION_NUMBER) &&\
pip install /code/onnxruntime/build/Linux/Release/dist/*.whl &&\
cd .. &&\
rm -rf onnxruntime cmake-3.14.3-Linux-x86_64

2
dockerfiles/README.md
View file

@ -88,6 +88,8 @@ Use `docker pull` with any of the images and tags below to pull an image and try
### **1. Using MCR container images**
*Note: ONNX Runtime 1.7 will be the last release that will publish MCR container images. Please switch to using PyPi packages or bulding from Dockefile section below from ONNX Runtime 1.8 onwards.*
The unified MCR container image can be used to run an application on any of the target accelerators. In order to select the target accelerator, the application should explicitly specifiy the choice using the *device_type* configuration option for OpenVINO Execution provider. Refer to [OpenVINO EP runtime configuration documentation](https://github.com/microsoft/onnxruntime/blob/master/docs/execution_providers/OpenVINO-ExecutionProvider.md#runtime-configuration-options) for details on specifying this option in the application code.
If the *device_type* runtime config option is not explicitly specified, CPU will be chosen as the hardware target execution.
### **2. Building from Dockerfile**

2
docs/Android_testing.md
View file

@ -1,6 +1,6 @@
# Testing Android Changes using the Emulator
See [BUILD.md](../Build.md#Android) for Android build instructions and information on the locations of the various files referred to here.
See [Android build instructions](https://www.onnxruntime.ai/docs/how-to/build.html#android) and information on the locations of the various files referred to here.
## Install the emulator

2
docs/ONNX_Runtime_Server_Usage.md
View file

@ -101,7 +101,7 @@ For easy tracking of requests, we provide the following header fields:
### rsyslog Support
If you prefer using an ONNX Runtime Server with [rsyslog](https://www.rsyslog.com/) support([build instruction](../BUILD.md#build-onnx-runtime-server-on-linux)), you should be able to see the log in `/var/log/syslog` after the ONNX Runtime Server runs. For detail about how to use rsyslog, please reference [here](https://www.rsyslog.com/category/guides-for-rsyslog/).
If you prefer using an ONNX Runtime Server with [rsyslog](https://www.rsyslog.com/) support([build instruction](https://www.onnxruntime.ai/docs/how-to/build.html#build-onnx-runtime-server-on-linux)), you should be able to see the log in `/var/log/syslog` after the ONNX Runtime Server runs. For detail about how to use rsyslog, please reference [here](https://www.rsyslog.com/category/guides-for-rsyslog/).
## Report Issues

44
docs/ONNX_Runtime_for_Mobile_Platforms.md
View file

@ -2,16 +2,16 @@
## Overview
<img align="left" width=40% src="images/Mobile.png" alt="Steps to build the reduced binary size."/>
<img align="left" width=40% src="images/Mobile.png" alt="Steps to build for mobile platforms."/>
ONNX Runtime now supports an internal model format to minimize the build size for usage in mobile and embedded scenarios. An ONNX model can be converted to an internal ONNX Runtime format ('ORT format model') using the below instructions.
The minimal build can be used with any ORT format model, provided that the kernels for the operators used in the model were included in the build.
i.e. the custom build provides a set of kernels, and if that set satisfies a given ORT format model's needs, the model can be loaded and executed.
A minimal build can be used with any ORT format model, provided that the kernels for the operators used in the model were included in the build.
I.e., the custom build provides a set of kernels, and if that set satisfies a given ORT format model's needs, the model can be loaded and executed.
## Steps to create model and minimal build
You will need a script from the the ONNX Runtime repository, and to also perform a custom build, so you will need to clone the repository locally. See [here](https://www.onnxruntime.ai/docs/how-to/build.html#prerequisites) for initial steps.
You will need a script from the ONNX Runtime repository and to also perform a custom build, so you will need to clone the repository locally. See [here](https://www.onnxruntime.ai/docs/how-to/build.html#prerequisites) for initial steps.
The directory the ONNX Runtime repository was cloned into is referred to as `<ONNX Runtime repository root>` in this documentation.
@ -20,21 +20,36 @@ Once you have cloned the repository, perform the following steps to create a min
### 1. Create ORT format model and configuration file with required operators
We will use a helper python script to convert ONNX format models into ORT format models, and to create the configuration file for use with the minimal build.
This will require the standard ONNX Runtime python package to be installed.
- Install the ONNX Runtime python package from https://pypi.org/project/onnxruntime/. Version 1.5.2 or later is required.
- `pip install onnxruntime`
- ensure that any existing ONNX Runtime python package was uninstalled first, or use `-U` with the above command to upgrade an existing package
- Copy all the ONNX models you wish to convert and use with the minimal build into a directory
- Convert the ONNX models to ORT format
- `python <ONNX Runtime repository root>/tools/python/convert_onnx_models_to_ort.py <path to directory containing one or more .onnx models>`
- For each ONNX model an ORT format model will be created with '.ort' as the file extension.
- A `required_operators.config` configuration file will also be created.
The configuration file specifies what operator kernels to include in the build.
This allows unused operator kernels to be pruned in order to decrease the binary size.
It is also possible (and optional) to further prune the operator kernel implementations based on their input and output type usage detected in the ORT format models.
This pruning is referred to as "operator type reduction" in this documentation.
- The helper python script requires the standard ONNX Runtime python package to be installed. Install the ONNX Runtime python package from https://pypi.org/project/onnxruntime/. Version 1.5.2 or later is required.
To enable operator type reduction, version 1.7 or later is required.
- `pip install onnxruntime`
- Ensure that any existing ONNX Runtime python package was uninstalled first, or use `-U` with the above command to upgrade an existing package.
- Additionally, if you want to enable operator type reduction, the Flatbuffers python package should be installed.
- `pip install flatbuffers`
- Copy all the ONNX models you wish to convert and use with the minimal build into a directory.
- Convert the ONNX models to ORT format
- `python <ONNX Runtime repository root>/tools/python/convert_onnx_models_to_ort.py <path to directory containing one or more .onnx models>`
- To enable operator type reduction, specify the `--enable_type_reduction` option.
- For each ONNX model an ORT format model will be created with '.ort' as the file extension.
- A configuration file will also be created.
If operator type reduction is enabled, the file will be called `required_operators_and_types.config`.
Otherwise, the file will be called `required_operators.config`.
Example:
Running `'python <ORT repository root>/tools/python/convert_onnx_models_to_ort.py /models'` where the '/models' directory contains ModelA.onnx and ModelB.onnx
- Will create /models/ModelA.ort and /models/ModelB.ort
- Will create /models/required_operators.config/
- Will create /models/required_operators.config
### 2. Create the minimal build
@ -48,6 +63,7 @@ See [here](https://www.onnxruntime.ai/docs/how-to/build.html#cpu) for the genera
The follow options can be used to reduce the build size. Enable all options that your scenario allows.
- Reduce build to required operator kernels
- Add `--include_ops_by_config <config file produced by step 1> --skip_tests` to the build parameters.
- To enable operator type reduction, also add `--enable_reduced_operator_type_support`.
- See the documentation on the [Reduced Operator Kernel build](Reduced_Operator_Kernel_build.md) for more information. This step can also be done pre-build if needed.
- NOTE: This step will edit some of the ONNX Runtime source files to exclude unused kernels. If you wish to go back to creating a full build, or wish to change the operator kernels included, you should run `git reset --hard` or `git checkout HEAD -- ./onnxruntime/core/providers` to undo these changes.

2
docs/Reduced_Operator_Kernel_build.md
View file

@ -6,7 +6,7 @@ A configuration file must be created with details of the kernels that are requir
Following that, ORT must be manually built, providing the configuration file in the `--include_ops_by_config` parameter. The build process will update the ORT kernel registration source files to exclude the unused kernels.
See the [build instructions](https://github.com/microsoft/onnxruntime/blob/master/BUILD.md#build-instructions) for more details on building ORT.
See the [build instructions](https://www.onnxruntime.ai/docs/how-to/build.html#build-instructions) for more details on building ORT.
When building ORT with a reduced set of kernel registrations, `--skip_tests` **MUST** be specified as the kernel reduction will render many of the unit tests invalid.

1
docs/Versioning.md
View file

@ -26,6 +26,7 @@ For more details on ONNX Release versions, see [this page](https://github.com/on
| ONNX Runtime release version | ONNX release version | ONNX opset version | ONNX ML opset version | Supported ONNX IR version | [Windows ML Availability](https://docs.microsoft.com/en-us/windows/ai/windows-ml/release-notes/)|
|------------------------------|--------------------|--------------------|----------------------|------------------|------------------|
| 1.7.0 | **1.8** down to 1.2 | 13 | 2 | 7 | Windows AI 1.7+ |
| 1.6.0 | **1.8** down to 1.2 | 13 | 2 | 7 | Windows AI 1.6+ |
| 1.5.3 | **1.7** down to 1.2 | 12 | 2 | 7 | Windows AI 1.5+ |
| 1.5.2 | **1.7** down to 1.2 | 12 | 2 | 7 | Windows AI 1.5+ |

5
docs/python/README.rst
View file

@ -8,6 +8,11 @@ For more information on ONNX Runtime, please see `aka.ms/onnxruntime <https://ak
Changes
-------
1.7.0
^^^^^
Release Notes : https://github.com/Microsoft/onnxruntime/releases/tag/v1.7.0
1.6.0
^^^^^

2
docs/python/notebooks/onnxruntime-nuphar-tutorial.ipynb
View file

@ -17,7 +17,7 @@
"\n",
"This example shows how to accelerate model inference using Nuphar, an execution provider that leverages just-in-time compilation to generate optimized executables.\n",
"\n",
"For more background about Nuphar, please check [Nuphar-ExecutionProvider.md](https://github.com/microsoft/onnxruntime/blob/master/docs/execution_providers/Nuphar-ExecutionProvider.md) and its [build instructions](https://github.com/microsoft/onnxruntime/blob/master/BUILD.md#nuphar).\n",
"For more background about Nuphar, please check [Nuphar-ExecutionProvider.md](https://github.com/microsoft/onnxruntime/blob/master/docs/execution_providers/Nuphar-ExecutionProvider.md) and its [build instructions](https://www.onnxruntime.ai/docs/how-to/build.html#nuphar).\n",
"\n",
"#### Tutorial Roadmap:\n",
"1. Prerequistes\n",

50
include/onnxruntime/core/common/make_string.h
View file

@ -19,10 +19,12 @@
#include <locale>
#include <sstream>
#include <type_traits>
namespace onnxruntime {
namespace detail {
inline void MakeStringImpl(std::ostringstream& /*ss*/) noexcept {
}
@ -36,6 +38,39 @@ inline void MakeStringImpl(std::ostringstream& ss, const T& t, const Args&... ar
MakeStringImpl(ss, t);
MakeStringImpl(ss, args...);
}
// see MakeString comments for explanation of why this is necessary
template <typename... Args>
inline std::string MakeStringImpl(const Args&... args) noexcept {
std::ostringstream ss;
MakeStringImpl(ss, args...);
return ss.str();
}
//
// Infrastructure to convert char[n] to char* to reduce binary size
//
// default is to leave the type as is
template <class T>
struct if_char_array_make_ptr {
using type = T;
};
// specialization that matches an array reference, which is what the char array from a string literal
// used in a call to MakeString will be.
// if the type is a char[n] array we 'decay' it to a char* so that the usages can be folded.
template <class T, size_t N>
struct if_char_array_make_ptr<T (&)[N]> {
// remove a single extent (T[x] -> T, but T[x][y] -> T[y]) so we only match char[x],
// and get the type name without the 'const' so both 'const char (&)[n]' and 'char (&)[n]' are matched.
using element_type = typename std::remove_const<typename std::remove_extent<T>::type>::type;
using type = typename std::conditional<std::is_same<char, element_type>::value, T*, T (&)[N]>::type;
};
// helper to make usage simpler in MakeString
template <class T>
using if_char_array_make_ptr_t = typename if_char_array_make_ptr<T>::type;
} // namespace detail
/**
@ -44,9 +79,18 @@ inline void MakeStringImpl(std::ostringstream& ss, const T& t, const Args&... ar
*/
template <typename... Args>
std::string MakeString(const Args&... args) {
std::ostringstream ss;
detail::MakeStringImpl(ss, args...);
return ss.str();
// We need to update the types from the MakeString template instantiation to decay any char[n] to char*.
// e.g. MakeString("in", "out") goes from MakeString<char[2], char[3]> to MakeStringImpl<char*, char*>
// so that MakeString("out", "in") will also match MakeStringImpl<char*, char*> instead of requiring
// MakeStringImpl<char[3], char[2]>.
//
// We have to do the type processing before any actual work, so this function purely implements the type processing.
// If we do not do it this way we do not get the full binary size reduction.
//
// See https://stackoverflow.com/a/29418212/684911 for overall details of the approach, but note it does not cover
// the need to do the type processing as a separate step.
return detail::MakeStringImpl(detail::if_char_array_make_ptr_t<Args const&>(args)...);
}
/**

11
include/onnxruntime/core/providers/coreml/coreml_provider_factory.h
View file

@ -19,17 +19,22 @@ enum COREMLFlags {
// Enable CoreML EP on subgraph
COREML_FLAG_ENABLE_ON_SUBGRAPH = 0x002,
// By default CoreML Execution provider will be enabled for all compatible Apple devices
// Enable this option will only enable CoreML EP for Apple devices with ANE (Apple Neural Engine)
// Please note, enable this option does not guarantee the entire model to be executed using ANE only
COREML_FLAG_ONLY_ENABLE_DEVICE_WITH_ANE = 0x004,
// Keep COREML_FLAG_MAX at the end of the enum definition
// And assign the last COREMLFlag to it
COREML_FLAG_LAST = COREML_FLAG_ENABLE_ON_SUBGRAPH,
COREML_FLAG_LAST = COREML_FLAG_ONLY_ENABLE_DEVICE_WITH_ANE,
};
#ifdef __cplusplus
extern "C" {
#endif
ORT_API_STATUS(OrtSessionOptionsAppendExecutionProvider_CoreML,
_In_ OrtSessionOptions* options, uint32_t coreml_flags);
ORT_EXPORT ORT_API_STATUS(OrtSessionOptionsAppendExecutionProvider_CoreML,
_In_ OrtSessionOptions* options, uint32_t coreml_flags);
#ifdef __cplusplus
}

1
include/onnxruntime/core/providers/cpu/cpu_provider_factory.h
View file

@ -10,6 +10,7 @@ extern "C" {
/**
* \param use_arena zero: false. non-zero: true.
*/
ORT_EXPORT
ORT_API_STATUS(OrtSessionOptionsAppendExecutionProvider_CPU, _In_ OrtSessionOptions* options, int use_arena)
ORT_ALL_ARGS_NONNULL;

3
include/onnxruntime/core/providers/nnapi/nnapi_provider_factory.h
View file

@ -44,7 +44,8 @@ enum NNAPIFlags {
extern "C" {
#endif
ORT_API_STATUS(OrtSessionOptionsAppendExecutionProvider_Nnapi, _In_ OrtSessionOptions* options, uint32_t nnapi_flags);
ORT_EXPORT ORT_API_STATUS(OrtSessionOptionsAppendExecutionProvider_Nnapi,
_In_ OrtSessionOptions* options, uint32_t nnapi_flags);
#ifdef __cplusplus
}

11
include/onnxruntime/core/session/onnxruntime_c_api.h
View file

@ -53,7 +53,12 @@ extern "C" {
#define ORT_MUST_USE_RESULT
#define ORTCHAR_T wchar_t
#else
// To make symbols visible on macOS/iOS
#ifdef __APPLE__
#define ORT_EXPORT __attribute__((visibility("default")))
#else
#define ORT_EXPORT
#endif
#define ORT_API_CALL
#define ORT_MUST_USE_RESULT __attribute__((warn_unused_result))
#define ORTCHAR_T char
@ -1131,7 +1136,7 @@ struct OrtApi {
ORT_API2_STATUS(SetGlobalDenormalAsZero, _Inout_ OrtThreadingOptions* tp_options);
/**
* Use this API to create the configuration of an arena that can eventually be used to define
* Use this API to create the configuration of an arena that can eventually be used to define
* an arena based allocator's behavior
* \param max_mem - use 0 to allow ORT to choose the default
* \param arena_extend_strategy - use -1 to allow ORT to choose the default, 0 = kNextPowerOfTwo, 1 = kSameAsRequested
@ -1151,8 +1156,8 @@ struct OrtApi {
* (doc_string field of the GraphProto message within the ModelProto message).
* If it doesn't exist, an empty string will be returned.
* \param model_metadata - an instance of OrtModelMetadata
* \param allocator - allocator used to allocate the string that will be returned back
* \param value - is set to a null terminated string allocated using 'allocator'.
* \param allocator - allocator used to allocate the string that will be returned back
* \param value - is set to a null terminated string allocated using 'allocator'.
The caller is responsible for freeing it.
*/
ORT_API2_STATUS(ModelMetadataGetGraphDescription, _In_ const OrtModelMetadata* model_metadata,

2
java/README.md
View file

@ -10,7 +10,7 @@ For general purpose usage of the publicly distributed API, please see the [gener
### Building
Use the main project's [build instructions](../BUILD.md) with the `--build_java` option.
Use the main project's [build instructions](https://www.onnxruntime.ai/docs/how-to/build.html) with the `--build_java` option.
#### Requirements

6
java/src/main/java/ai/onnxruntime/OrtProvider.java
View file

@ -1,5 +1,5 @@
/*
* Copyright (c) 2020, Oracle and/or its affiliates. All rights reserved.
* Copyright (c) 2020, 2021, Oracle and/or its affiliates. All rights reserved.
* Licensed under the MIT License.
*/
package ai.onnxruntime;
@ -21,7 +21,9 @@ public enum OrtProvider {
DIRECT_ML("DmlExecutionProvider"),
MI_GRAPH_X("MIGraphXExecutionProvider"),
ACL("ACLExecutionProvider"),
ARM_NN("ArmNNExecutionProvider");
ARM_NN("ArmNNExecutionProvider"),
ROCM("ROCMExecutionProvider"),
CORE_ML("CoreMLExecutionProvider");
private static final Map<String, OrtProvider> valueMap = new HashMap<>(values().length);

70
java/src/main/java/ai/onnxruntime/OrtSession.java
View file

@ -4,10 +4,14 @@
*/
package ai.onnxruntime;
import ai.onnxruntime.providers.CoreMLFlags;
import ai.onnxruntime.providers.NNAPIFlags;
import ai.onnxruntime.providers.OrtFlags;
import java.io.IOException;
import java.util.ArrayList;
import java.util.Arrays;
import java.util.Collections;
import java.util.EnumSet;
import java.util.Iterator;
import java.util.LinkedHashMap;
import java.util.LinkedHashSet;
@ -809,13 +813,23 @@ public class OrtSession implements AutoCloseable {
}
/**
* Adds Android's NNAPI as an execution backend.
* Adds Android's NNAPI as an execution backend. Uses the default empty flag.
*
* @throws OrtException If there was an error in native code.
*/
public void addNnapi() throws OrtException {
addNnapi(EnumSet.noneOf(NNAPIFlags.class));
}
/**
* Adds Android's NNAPI as an execution backend.
*
* @param flags The flags which control the NNAPI configuration.
* @throws OrtException If there was an error in native code.
*/
public void addNnapi(EnumSet<NNAPIFlags> flags) throws OrtException {
checkClosed();
addNnapi(OnnxRuntime.ortApiHandle, nativeHandle, 0);
addNnapi(OnnxRuntime.ortApiHandle, nativeHandle, OrtFlags.aggregateToInt(flags));
}
/**
@ -852,6 +866,49 @@ public class OrtSession implements AutoCloseable {
addACL(OnnxRuntime.ortApiHandle, nativeHandle, useArena ? 1 : 0);
}
/**
* Adds the ARM Neural Net library as an execution backend.
*
* @param useArena If true use the arena memory allocator.
* @throws OrtException If there was an error in native code.
*/
public void addArmNN(boolean useArena) throws OrtException {
checkClosed();
addArmNN(OnnxRuntime.ortApiHandle, nativeHandle, useArena ? 1 : 0);
}
/**
* Adds ROCM as an execution backend.
*
* @param deviceID The ROCM device ID.
* @param memLimit The maximum amount of memory available.
* @throws OrtException If there was an error in native code.
*/
public void addROCM(int deviceID, long memLimit) throws OrtException {
checkClosed();
addROCM(OnnxRuntime.ortApiHandle, nativeHandle, deviceID, memLimit);
}
/**
* Adds Apple's CoreML as an execution backend. Uses the default empty flag.
*
* @throws OrtException If there was an error in native code.
*/
public void addCoreML() throws OrtException {
addCoreML(EnumSet.noneOf(CoreMLFlags.class));
}
/**
* Adds Apple's CoreML as an execution backend.
*
* @param flags The flags which control the CoreML configuration.
* @throws OrtException If there was an error in native code.
*/
public void addCoreML(EnumSet<CoreMLFlags> flags) throws OrtException {
checkClosed();
addCoreML(OnnxRuntime.ortApiHandle, nativeHandle, OrtFlags.aggregateToInt(flags));
}
private native void setExecutionMode(long apiHandle, long nativeHandle, int mode)
throws OrtException;
@ -944,6 +1001,15 @@ public class OrtSession implements AutoCloseable {
throws OrtException;
private native void addACL(long apiHandle, long nativeHandle, int useArena) throws OrtException;
private native void addArmNN(long apiHandle, long nativeHandle, int useArena)
throws OrtException;
private native void addROCM(long apiHandle, long nativeHandle, int deviceID, long memLimit)
throws OrtException;
private native void addCoreML(long apiHandle, long nativeHandle, int coreMLFlags)
throws OrtException;
}
/** Used to control logging and termination of a call to {@link OrtSession#run}. */

23
java/src/main/java/ai/onnxruntime/providers/CoreMLFlags.java Normal file
View file

@ -0,0 +1,23 @@
/*
* Copyright (c) 2021, Oracle and/or its affiliates. All rights reserved.
* Licensed under the MIT License.
*/
package ai.onnxruntime.providers;
/** Flags for the CoreML provider. */
public enum CoreMLFlags implements OrtFlags {
CPU_ONLY(1), // COREML_FLAG_USE_CPU_ONLY(0x001)
ENABLE_ON_SUBGRAPH(2), // COREML_FLAG_ENABLE_ON_SUBGRAPH(0x002)
ONLY_ENABLE_DEVICE_WITH_ANE(4); // COREML_FLAG_ONLY_ENABLE_DEVICE_WITH_ANE(0x004),
public final int value;
CoreMLFlags(int value) {
this.value = value;
}
@Override
public int getValue() {
return value;
}
}

23
java/src/main/java/ai/onnxruntime/providers/NNAPIFlags.java Normal file
View file

@ -0,0 +1,23 @@
/*
* Copyright (c) 2021, Oracle and/or its affiliates. All rights reserved.
* Licensed under the MIT License.
*/
package ai.onnxruntime.providers;
/** Flags for the NNAPI provider. */
public enum NNAPIFlags implements OrtFlags {
USE_FP16(1), // NNAPI_FLAG_USE_FP16(0x001)
USE_NCHW(2), // NNAPI_FLAG_USE_NCHW(0x002)
CPU_DISABLED(4); // NNAPI_FLAG_CPU_DISABLED(0x004)
public final int value;
NNAPIFlags(int value) {
this.value = value;
}
@Override
public int getValue() {
return value;
}
}

35
java/src/main/java/ai/onnxruntime/providers/OrtFlags.java Normal file
View file

@ -0,0 +1,35 @@
/*
* Copyright (c) 2021, Oracle and/or its affiliates. All rights reserved.
* Licensed under the MIT License.
*/
package ai.onnxruntime.providers;
import java.util.EnumSet;
/** An interface for bitset enums that should be aggregated into a single integer. */
public interface OrtFlags {
/**
* Gets the underlying flag value.
*
* @return The flag value.
*/
public int getValue();
/**
* Converts an EnumSet of flags into the value expected by the C API.
*
* @param set The enum set to aggregate the values from.
* @param <E> The enum type to aggregate.
* @return The aggregated values
*/
public static <E extends Enum<E> & OrtFlags> int aggregateToInt(EnumSet<E> set) {
int value = 0;
for (OrtFlags flag : set) {
value |= flag.getValue();
}
return value;
}
}

52
java/src/main/native/ai_onnxruntime_OrtSession_SessionOptions.c
View file

@ -23,6 +23,9 @@
#include "onnxruntime/core/providers/tensorrt/tensorrt_provider_factory.h"
#include "onnxruntime/core/providers/migraphx/migraphx_provider_factory.h"
#include "onnxruntime/core/providers/acl/acl_provider_factory.h"
#include "onnxruntime/core/providers/armnn/armnn_provider_factory.h"
#include "onnxruntime/core/providers/coreml/coreml_provider_factory.h"
#include "onnxruntime/core/providers/rocm/rocm_provider_factory.h"
#ifdef USE_DIRECTML
#include "onnxruntime/core/providers/dml/dml_provider_factory.h"
#endif
@ -502,3 +505,52 @@ JNIEXPORT void JNICALL Java_ai_onnxruntime_OrtSession_00024SessionOptions_addACL
throwOrtException(jniEnv,convertErrorCode(ORT_INVALID_ARGUMENT),"This binary was not compiled with ACL support.");
#endif
}
/*
* Class: ai_onnxruntime_OrtSession_SessionOptions
* Method: addArmNN
* Signature: (JJI)V
*/
JNIEXPORT void JNICALL Java_ai_onnxruntime_OrtSession_00024SessionOptions_addArmNN
(JNIEnv * jniEnv, jobject jobj, jlong apiHandle, jlong handle, jint useArena) {
(void)jobj;
#ifdef USE_ARMNN
checkOrtStatus(jniEnv,(const OrtApi*)apiHandle,OrtSessionOptionsAppendExecutionProvider_ArmNN((OrtSessionOptions*) handle,useArena));
#else
(void)apiHandle;(void)handle;(void)useArena; // Parameters used when ARMNN is defined.
throwOrtException(jniEnv,convertErrorCode(ORT_INVALID_ARGUMENT),"This binary was not compiled with ArmNN support.");
#endif
}
/*
* Class: ai_onnxruntime_OrtSession_SessionOptions
* Method: addCoreML
* Signature: (JJI)V
*/
JNIEXPORT void JNICALL Java_ai_onnxruntime_OrtSession_00024SessionOptions_addCoreML
(JNIEnv * jniEnv, jobject jobj, jlong apiHandle, jlong handle, jint coreMLFlags) {
(void)jobj;
#ifdef USE_CORE_ML
checkOrtStatus(jniEnv,(const OrtApi*)apiHandle,OrtSessionOptionsAppendExecutionProvider_CoreML((OrtSessionOptions*) handle, (uint32_t) coreMLFlags));
#else
(void)apiHandle;(void)handle;(void)coreMLFlags; // Parameters used when CoreML is defined.
throwOrtException(jniEnv,convertErrorCode(ORT_INVALID_ARGUMENT),"This binary was not compiled with CoreML support.");
#endif
}
/*
* Class: ai_onnxruntime_OrtSession_SessionOptions
* Method: addROCM
* Signature: (JJI)V
*/
JNIEXPORT void JNICALL Java_ai_onnxruntime_OrtSession_00024SessionOptions_addROCM
(JNIEnv * jniEnv, jobject jobj, jlong apiHandle, jlong handle, jint deviceID, jlong memLimit) {
(void)jobj;
#ifdef USE_ROCM
checkOrtStatus(jniEnv,(const OrtApi*)apiHandle,OrtSessionOptionsAppendExecutionProvider_ROCM((OrtSessionOptions*) handle, deviceID, (size_t) memLimit));
#else
(void)apiHandle;(void)handle;(void)deviceID;(void)memLimit; // Parameters used when ROCM is defined.
throwOrtException(jniEnv,convertErrorCode(ORT_INVALID_ARGUMENT),"This binary was not compiled with ROCM support.");
#endif
}

4
nodejs/README.md
View file

@ -16,7 +16,7 @@ Install the latest dev version:
npm install onnxruntime@dev
```
Refer to [Node.js samples](../samples/README.md#Nodejs) for samples and tutorials.
Refer to [Node.js samples](../samples/nodejs/README.md) for samples and tutorials.
## Requirements
@ -28,7 +28,7 @@ Following platforms are supported with pre-built binaries:
- Linux x64 CPU NAPI_v3
- MacOS x64 CPU NAPI_v3
To use on platforms without pre-built binaries, you can build Node.js binding from source and consume it by `npm install <onnxruntime_repo_root>/nodejs/`. See also [BUILD.MD](../BUILD.md#apis-and-language-bindings) for building ONNX Runtime Node.js binding locally.
To use on platforms without pre-built binaries, you can build Node.js binding from source and consume it by `npm install <onnxruntime_repo_root>/nodejs/`. See also [instructions](https://www.onnxruntime.ai/docs/how-to/build.html#apis-and-language-bindings) for building ONNX Runtime Node.js binding locally.
## License

19000
nodejs/package-lock.json generated
View file

File diff suppressed because it is too large Load diff

130
nodejs/package.json
View file

@ -1,72 +1,72 @@
{
"name": "onnxruntime",
"description": "Node.js binding of ONNXRuntime",
"version": "1.6.0",
"main": "./lib/index.js",
"types": "./types/lib/index.d.ts",
"scripts": {
"install": "prebuild-install -r napi || (tsc && node ./script/build)",
"build": "tsc && node ./script/build",
"buildd": "tsc && node ./script/build --config=Debug",
"buildr": "tsc && node ./script/build --config=RelWithDebInfo",
"rebuild": "tsc && node ./script/build --rebuild",
"rebuildd": "tsc && node ./script/build --rebuild --config=Debug",
"rebuildr": "tsc && node ./script/build --rebuild --config=RelWithDebInfo",
"test": "mocha ./test/test-main",
"lint": "eslint . --ext .ts",
"prepack": "node ./script/pack-prebuild",
"format": "clang-format --glob=\"{{lib,test,script}/**/*.ts,src/**/*.{cc,h}}\" --style=file -i"
},
"repository": {
"type": "git",
"url": "https://github.com/Microsoft/onnxruntime.git"
},
"keywords": [
"ONNX",
"ONNX Runtime"
],
"os": [
"win32",
"darwin",
"linux"
],
"binary": {
"module_path": "./bin",
"host": "https://onnxruntimetestdata.blob.core.windows.net/onnxruntime-node-prebuild/",
"module_path": "./bin",
"host": "https://onnxruntimetestdata.blob.core.windows.net/onnxruntime-node-prebuild/",
"napi_versions": [
3
]
},
"author": "fs-eire",
"license": "MIT",
"devDependencies": {
"@types/fs-extra": "^9.0.6",
"@types/klaw-sync": "^6.0.0",
"@types/minimist": "1.2.1",
"@types/mocha": "^8.2.0",
"@types/tar-stream": "^2.2.0",
"@typescript-eslint/eslint-plugin": "^4.14.2",
"@typescript-eslint/parser": "^4.14.2",
"clang-format": "^1.5.0",
"cmake-js": "^6.1.0",
"eslint": "^7.19.0",
"eslint-plugin-import": "^2.22.1",
"eslint-plugin-jsdoc": "^31.6.0",
"eslint-plugin-prefer-arrow": "^1.2.3",
"fs-extra": "^9.1.0",
"globby": "^11.0.2",
"jsonc": "^2.0.0",
"klaw-sync": "^6.0.0",
"minimist": "^1.2.5",
"mocha": "^8.2.1",
"node-addon-api": "^3.1.0",
"node-pre-gyp-github": "^1.4.3",
"onnx-proto": "^4.0.4",
"tar-stream": "^2.2.0",
"typedoc": "^0.20.25",
"typescript": "^4.1.3"
},
},
"license": "MIT",
"name": "onnxruntime",
"repository": {
"url": "https://github.com/Microsoft/onnxruntime.git",
"type": "git"
},
"author": "fs-eire",
"version": "1.7.0",
"dependencies": {
"prebuild-install": "^6.0.0"
}
}
},
"scripts": {
"lint": "eslint . --ext .ts",
"buildr": "tsc && node ./script/build --config=RelWithDebInfo",
"format": "clang-format --glob=\"{{lib,test,script}/**/*.ts,src/**/*.{cc,h}}\" --style=file -i",
"rebuild": "tsc && node ./script/build --rebuild",
"rebuildd": "tsc && node ./script/build --rebuild --config=Debug",
"buildd": "tsc && node ./script/build --config=Debug",
"build": "tsc && node ./script/build",
"install": "prebuild-install -r napi || (tsc && node ./script/build)",
"test": "mocha ./test/test-main",
"prepack": "node ./script/pack-prebuild",
"rebuildr": "tsc && node ./script/build --rebuild --config=RelWithDebInfo"
},
"keywords": [
"ONNX",
"ONNX Runtime"
],
"devDependencies": {
"typedoc": "^0.20.25",
"mocha": "^8.2.1",
"@types/fs-extra": "^9.0.6",
"@types/tar-stream": "^2.2.0",
"klaw-sync": "^6.0.0",
"minimist": "^1.2.5",
"@typescript-eslint/parser": "^4.14.2",
"clang-format": "^1.5.0",
"@types/klaw-sync": "^6.0.0",
"node-addon-api": "^3.1.0",
"cmake-js": "^6.1.0",
"eslint-plugin-prefer-arrow": "^1.2.3",
"typescript": "^4.1.3",
"jsonc": "^2.0.0",
"@types/mocha": "^8.2.0",
"node-pre-gyp-github": "^1.4.3",
"@typescript-eslint/eslint-plugin": "^4.14.2",
"eslint-plugin-jsdoc": "^31.6.0",
"onnx-proto": "^4.0.4",
"globby": "^11.0.2",
"fs-extra": "^9.1.0",
"eslint": "^7.19.0",
"tar-stream": "^2.2.0",
"@types/minimist": "1.2.1",
"eslint-plugin-import": "^2.22.1"
},
"main": "./lib/index.js",
"os": [
"win32",
"darwin",
"linux"
],
"types": "./types/lib/index.d.ts",
"description": "Node.js binding of ONNXRuntime"
}

13
onnxruntime/__init__.py
View file

@ -7,20 +7,9 @@ ONNX Runtime is a performance-focused scoring engine for Open Neural Network Exc
For more information on ONNX Runtime, please see `aka.ms/onnxruntime <https://aka.ms/onnxruntime/>`_
or the `Github project <https://github.com/microsoft/onnxruntime/>`_.
"""
__version__ = "1.6.0"
__version__ = "1.7.0"
__author__ = "Microsoft"
import os
import platform
import sys
# Python 3.8 (and later) on Windows doesn't search system PATH when loading DLLs,
# so CUDA location needs to be specified explicitly. This needs to be done before importing
# onnxruntime.capi._pybind_state
if "CUDA_PATH" in os.environ and platform.system() == "Windows" and sys.version_info >= (3, 8):
cuda_bin_dir = os.path.join(os.environ["CUDA_PATH"], "bin")
os.add_dll_directory(cuda_bin_dir)
from onnxruntime.capi._pybind_state import get_all_providers, get_available_providers, get_device, set_seed, \
RunOptions, SessionOptions, set_default_logger_severity, enable_telemetry_events, disable_telemetry_events, \
NodeArg, ModelMetadata, GraphOptimizationLevel, ExecutionMode, ExecutionOrder, OrtDevice, SessionIOBinding, \

5
onnxruntime/contrib_ops/cpu/bert/longformer_attention_base.h
View file

@ -25,5 +25,10 @@ class LongformerAttentionBase {
int window_; // Attention windows length (W). It is half (one-sided) of total window size.
};
namespace longformer {
// Environment variable to give a hint about choosing kernels for less memory or latency.
constexpr const char* kUseCompactMemory = "ORT_LONGFORMER_COMPACT_MEMORY";
} // namespace longformer
} // namespace contrib
} // namespace onnxruntime

20
onnxruntime/contrib_ops/cuda/bert/longformer_attention.cc
View file

@ -5,6 +5,7 @@
#include "core/framework/tensorprotoutils.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/shared_inc/fpgeneric.h"
#include "core/platform/env_var_utils.h"
#include "longformer_global_impl.h"
#include "longformer_attention_impl.h"
@ -49,7 +50,9 @@ class AutoDestoryCudaEvent {
};
template <typename T>
LongformerAttention<T>::LongformerAttention(const OpKernelInfo& info) : CudaKernel(info), LongformerAttentionBase(info) {}
LongformerAttention<T>::LongformerAttention(const OpKernelInfo& info) : CudaKernel(info), LongformerAttentionBase(info) {
use_compact_memory_ = ParseEnvironmentVariableWithDefault<bool>(longformer::kUseCompactMemory, false);
}
template <typename T>
Status LongformerAttention<T>::ComputeInternal(OpKernelContext* context) const {
@ -80,6 +83,7 @@ Status LongformerAttention<T>::ComputeInternal(OpKernelContext* context) const {
constexpr size_t element_size = sizeof(T);
// TODO: only calculate once per model.
// Build Global Index
auto global_index_buffer = GetScratchBuffer<int>(batch_size * sequence_length);
auto batch_global_num_buffer = GetScratchBuffer<int>(batch_size);
@ -148,10 +152,11 @@ Status LongformerAttention<T>::ComputeInternal(OpKernelContext* context) const {
}
}
// Cuda kernel implementation has a limitation of number of global tokens.
if (max_num_global > window_) {
ORT_THROW("LongformerAttention CUDA operator does not support number of global tokens > attention window.");
}
// Force to use fast kernel in two situations:
// (1) global tokens > windows size. In that case, compact memory kernel cannot be used.
// (2) sequence_length == 2 * attention_window. Use fast kernel to walk around parity issue of compact memory kernel.
// In other case, we will choose according to user's environment variable setting (default is fast kernel).
bool use_fast_kernel = (max_num_global > window_ || sequence_length == 2 * window_ || !use_compact_memory_);
// Fully connection for global projection.
// Note that Q only need handle global query tokens if we split GEMM to global Q/K/V separately.
@ -172,7 +177,7 @@ Status LongformerAttention<T>::ComputeInternal(OpKernelContext* context) const {
&one, reinterpret_cast<CudaT*>(global_gemm_buffer.get()), n, device_prop));
}
size_t workSpaceSize = GetLongformerAttentionWorkspaceSize(element_size, batch_size, num_heads_, head_size, sequence_length, max_num_global, window_);
size_t workSpaceSize = GetLongformerAttentionWorkspaceSize(element_size, batch_size, num_heads_, head_size, sequence_length, max_num_global, window_, use_fast_kernel);
auto workspace_buffer = GetScratchBuffer<void>(workSpaceSize);
if (!LaunchLongformerAttentionKernel(
device_prop,
@ -193,7 +198,8 @@ Status LongformerAttention<T>::ComputeInternal(OpKernelContext* context) const {
head_size,
window_,
max_num_global,
element_size)) {
element_size,
use_fast_kernel)) {
// Get last error to reset it to cudaSuccess.
CUDA_CALL(cudaGetLastError());
return Status(common::ONNXRUNTIME, common::FAIL);

3
onnxruntime/contrib_ops/cuda/bert/longformer_attention.h
View file

@ -18,6 +18,9 @@ class LongformerAttention final : public CudaKernel, public LongformerAttentionB
public:
LongformerAttention(const OpKernelInfo& info);
Status ComputeInternal(OpKernelContext* context) const override;
private:
bool use_compact_memory_;
};
} // namespace cuda

135
onnxruntime/contrib_ops/cuda/bert/longformer_attention_impl.cu
View file

@ -28,7 +28,7 @@ limitations under the License.
#include "core/providers/cuda/cuda_common.h"
#include "longformer_attention_impl.h"
#include "attention_impl.h"
#include "attention_softmax.h"
#include "longformer_attention_softmax.h"
using namespace onnxruntime::cuda;
using namespace cub;
@ -53,11 +53,8 @@ namespace cuda {
// [SoftmaxSpace: see below] [Q:BxNxSxH] [K:BxNxSxH] [V:BxNxSxH] [Global_Q:BxNxSxH] [Global_K:BxNxSxH] [Global_V:BxNxSxH]
// where Global_Q, Global_K and Global_V are optional. They are not allocated when there is no global token.
//
// It is feasible to use compact format for Global_Q with shape BxNxGxH to save space. We do not use compact format for now.
//
// SoftmaxSpace layout:
// [scratch1: (5S-3W)*W*N*B][scratch2: size_t 20]
//
// Scratch1 has 5 buffers for local and global attention calculation.
// Scratch2 has 5 input pointers, 5 output pointers, 5 buffer sizes and 5 strides related to scratch1.
@ -74,10 +71,17 @@ size_t GetLongformerSoftmaxWorkspaceSize(
int batch_size,
int num_heads,
int sequence_length,
int window) {
size_t scratch1_size = GetScratch1Size(element_size, batch_size, num_heads, sequence_length, window);
size_t scratch2_size = 10 * (sizeof(void*) + sizeof(size_t));
return scratch1_size + scratch2_size;
int window,
bool use_fast_kernel) {
if (!use_fast_kernel) {
size_t scratch1_size = GetScratch1Size(element_size, batch_size, num_heads, sequence_length, window);
size_t scratch2_size = 10 * (sizeof(void*) + sizeof(size_t));
return scratch1_size + scratch2_size;
} else {
// Non-compact layout when environment variable ORT_LONGFORMER_COMPACT_MEMORY=0 is set.
// [scratch1: BxNxSxS] [scratch2: BxNxSxS]
return 2 * GetAttentionScratchSize(element_size, batch_size, num_heads, sequence_length, sequence_length);
}
}
size_t GetLongformerAttentionWorkspaceSize(
@ -87,8 +91,9 @@ size_t GetLongformerAttentionWorkspaceSize(
int head_size,
int sequence_length,
int max_num_global,
int window) {
size_t softmax_size = GetLongformerSoftmaxWorkspaceSize(element_size, batch_size, num_heads, sequence_length, window);
int window,
bool use_fast_kernel) {
size_t softmax_size = GetLongformerSoftmaxWorkspaceSize(element_size, batch_size, num_heads, sequence_length, window, use_fast_kernel);
size_t qkv_size = 3 * batch_size * sequence_length * num_heads * head_size * element_size;
size_t global_qkv_size = max_num_global > 0 ? qkv_size : 0;
return softmax_size + qkv_size + global_qkv_size;
@ -100,6 +105,7 @@ size_t GetPinnedBufferSize(int batch_size) {
return sizeof(int) * batch_size + GetScratch2Size();
}
// Softmax kernel for compact format
template <typename T, int blockSize>
__launch_bounds__(blockSize)
__global__ void LongformerSoftmaxKernel(const int* global_attention,
@ -354,7 +360,6 @@ bool launchSoftmaxKernel(
cudaStream_t stream,
cublasHandle_t cublas,
void* workspace,
size_t softmax_workspace_size,
const void* q, // transposed Q with shape (B, N, S, H)
const void* k, // transposed K with shape (B, N, S, H)
const void* v, // transposed V with shape (B, N, S, H)
@ -373,10 +378,7 @@ bool launchSoftmaxKernel(
int num_heads, // number of heads
int head_size, // hidden size per head
int window, // one sided window size
int max_num_global, // maximum number of global tokens (G) in all batches
size_t element_size) { // size of element: 2 for half, and 4 for float
assert(max_num_global <= window);
const int* global_count = reinterpret_cast<const int*>(pinned_buffer);
bool is_fp16 = (element_size == 2);
@ -605,7 +607,10 @@ bool launchSoftmaxKernel(
resultType,
algo));
void* global_q_batch = (char*)global_q + (i * elements_per_batch) * element_size; // For compact format: replace elements_per_batch by num_heads * max_num_global * head_size
// It is feasible to use compact format for Global_Q with shape BxNxGxH to save space.
// In that case, elements_per_batch is num_heads * max_num_global * head_size, and stride_per_head is max_num_global * head_size.
void* global_q_batch = (char*)global_q + (i * elements_per_batch) * element_size;
void* global_k_batch = (char*)global_k + (i * elements_per_batch) * element_size;
qk_batch = (char*)input_pointers[4] + (i * buffer_sizes[4] * num_heads) * element_size;
@ -625,7 +630,7 @@ bool launchSoftmaxKernel(
global_q_batch, // B
Btype, // B type
head_size, // ldb
stride_per_head, // strideB. For compact format: max_num_global * head_size.
stride_per_head, // strideB.
beta_0, // beta
qk_batch, // C
Ctype, // C type
@ -827,8 +832,9 @@ bool LongformerQkvToContext(
const T* global_input, const int* global_attention,
const int* global_index, const int* batch_global_num, const int max_num_global,
void* pinned_buffer, T* workspace,
T* output) {
size_t softmax_workspace_size = GetLongformerSoftmaxWorkspaceSize(element_size, batch_size, num_heads, sequence_length, window);
T* output,
size_t softmax_workspace_size,
bool use_fast_kernel) {
T* qkv = reinterpret_cast<T*>((char*)workspace + softmax_workspace_size);
// Number of elements in Q, K, V, Global_Q, Global_K or Global_V are same: BxNxSxH
@ -862,33 +868,62 @@ bool LongformerQkvToContext(
const float rsqrt_head_size = 1.f / sqrt(static_cast<float>(head_size));
T* temp_output = qkv; // Q will be overwritten
if (!launchSoftmaxKernel(
stream,
cublas,
workspace,
softmax_workspace_size,
q, // Transposed Q with shape B x N x S x H
k, // Transposed K with shape B x N x S x H
v, // Transposed V with shape B x N x S x H
attention_mask, // Attention mask flags with shape B x S
global_q, // Transposed global Q with shape B x N x S x H.
global_k, // Transposed global K with shape B x N x S x H
global_v, // Transposed global V with shape B x N x S x H
global_attention, // Global attention flags with shape B x S
global_index, // Global index with shape B x S
batch_global_num, // Number of global token per batch with shape B x 1
pinned_buffer, // Pinned Memory Buffer
temp_output, // Output with shape B x N x S x H
rsqrt_head_size, // Scaler
batch_size, // Batch size
sequence_length, // Sequence length
num_heads, // Number of attention heads
head_size, // Hidden size per head
window, // Half (one-sided) window size
max_num_global, // Maximum number of global tokens (G)
element_size)) {
return false;
if (use_fast_kernel) {
if (!launchSoftmaxFastKernel(
stream,
cublas,
workspace, // softmax space
q, // transposed Q with shape (B, N, S, H)
k, // transposed K with shape (B, N, S, H)
v, // transposed V with shape (B, N, S, H)
attention_mask, // attention mask with shape (B, S), with value 0.0 not masked, and -10000.0 masked.
global_q, // Q for global tokens with shape (B, N, S, H)
global_k, // K for global tokens with shape (B, N, S, H)
global_v, // V for global tokens with shape (B, N, S, H)
global_attention, // global attention with shape (B, S), with value 0 for local attention and 1 for global attention.
global_index, // Global index with shape (B, S)
batch_global_num, // Number of global tokens per batch with shape (B, 1)
pinned_buffer, // Pinned memory in CPU. Number of global tokens per batch with shape (B, 1)
temp_output, // output with shape (B, N, S, H)
rsqrt_head_size, // scalar
batch_size, // batch size
sequence_length, // sequence length
num_heads, // number of heads
head_size, // hidden size per head
window, // Half (one-sided) window size
element_size)) {
return false;
}
} else {
assert(max_num_global <= window);
if (!launchSoftmaxKernel(
stream,
cublas,
workspace, // softmax space
q, // Transposed Q with shape B x N x S x H
k, // Transposed K with shape B x N x S x H
v, // Transposed V with shape B x N x S x H
attention_mask, // Attention mask flags with shape B x S. Value -10000.0 means masked, and 0.0 not mased.
global_q, // Transposed global Q with shape B x N x S x H.
global_k, // Transposed global K with shape B x N x S x H
global_v, // Transposed global V with shape B x N x S x H
global_attention, // Global attention flags with shape B x S
global_index, // Global index with shape B x S
batch_global_num, // Number of global token per batch with shape B x 1
pinned_buffer, // Pinned Memory Buffer
temp_output, // Output with shape B x N x S x H
rsqrt_head_size, // Scaler
batch_size, // Batch size
sequence_length, // Sequence length
num_heads, // Number of attention heads
head_size, // Hidden size per head
window, // Half (one-sided) window size
element_size)) {
return false;
}
}
// The temp_output is BxNxSxH, transpose it to final output BxSxNxH
return LaunchTransCtx(stream, sequence_length, batch_size, head_size, num_heads, temp_output, output);
@ -913,8 +948,10 @@ bool LaunchLongformerAttentionKernel(
int head_size,
int window,
int max_num_global,
const size_t element_size) {
const size_t element_size,
bool use_fast_kernel) {
CublasMathModeSetter helper(device_prop, cublas, CUBLAS_TENSOR_OP_MATH);
size_t softmax_workspace_size = GetLongformerSoftmaxWorkspaceSize(element_size, batch_size, num_heads, sequence_length, window, use_fast_kernel);
if (element_size == 2) {
return LongformerQkvToContext(cublas, stream,
batch_size, sequence_length, num_heads, head_size, window, element_size,
@ -927,7 +964,9 @@ bool LaunchLongformerAttentionKernel(
max_num_global,
pinned_buffer,
reinterpret_cast<half*>(workspace),
reinterpret_cast<half*>(output));
reinterpret_cast<half*>(output),
softmax_workspace_size,
use_fast_kernel);
} else {
return LongformerQkvToContext(cublas, stream,
batch_size, sequence_length, num_heads, head_size, window, element_size,
@ -940,7 +979,9 @@ bool LaunchLongformerAttentionKernel(
max_num_global,
pinned_buffer,
reinterpret_cast<float*>(workspace),
reinterpret_cast<float*>(output));
reinterpret_cast<float*>(output),
softmax_workspace_size,
use_fast_kernel);
}
}

8
onnxruntime/contrib_ops/cuda/bert/longformer_attention_impl.h
View file

@ -18,9 +18,10 @@ size_t GetLongformerAttentionWorkspaceSize(
int head_size,
int sequence_length,
int max_num_global,
int window);
int window,
bool use_fast_kernel);
bool LaunchLongformerAttentionKernel(
bool LaunchLongformerAttentionKernel(
const cudaDeviceProp& device_prop, // Device Properties
cublasHandle_t& cublas, // Cublas handle
cudaStream_t stream, // CUDA stream
@ -39,7 +40,8 @@ size_t GetLongformerAttentionWorkspaceSize(
int head_size, // Hidden layer size per head (H)
int window, // One sided attention window (W)
int max_num_global, // Maximum number of global tokens (G)
const size_t element_size // Element size of input tensor
const size_t element_size, // Element size of input tensor,
bool use_fast_kernel // Use compact memory
);
} // namespace cuda

657
onnxruntime/contrib_ops/cuda/bert/longformer_attention_softmax.cu Normal file
View file

@ -0,0 +1,657 @@
/*
Copyright (c) NVIDIA Corporation and Microsoft Corporation
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 is fast cuda kernels for longformer attention softmax.
// It uses two temporary matrix of BxNxSxS, and consumes more memory when sequence length is large.
#include <cub/cub.cuh>
#include <cublas_v2.h>
#include <cuda_fp16.h>
#include <cuda_runtime.h>
#include <math_constants.h>
#include "core/providers/cuda/cu_inc/common.cuh"
#include "core/providers/cuda/cuda_common.h"
#include "longformer_attention_softmax.h"
#include "attention_impl.h"
using namespace onnxruntime::cuda;
using namespace cub;
#define CHECK(expr) \
if (!CUBLAS_CALL(expr)) { \
return false; \
}
namespace onnxruntime {
namespace contrib {
namespace cuda {
template <typename T, int blockSize>
__launch_bounds__(blockSize)
__global__ void LongformerSoftmaxFastKernel(const int* global_attention,
const int* global_index,
const int* batch_global_num,
const T* input,
const T* attention_mask,
T* output,
float scaler,
int dim0,
int sequence_length,
int attention_window) {
typedef cub::BlockReduce<float, blockSize> BlockReduce;
__shared__ typename BlockReduce::TempStorage block_reduce_temp;
__shared__ float max_shared;
__shared__ float sum_shared;
const T* input_block = input + sequence_length * blockIdx.x;
T* output_block = output + sequence_length * blockIdx.x;
const int batch_index = blockIdx.x / dim0;
const int row_index = blockIdx.x % sequence_length;
const int global_num = batch_global_num[batch_index];
// To be consistent with Huggingface Longformer, the row of maksed word are set as zero.
if ((float)attention_mask[batch_index * sequence_length + row_index] < 0.0f) {
for (int i = threadIdx.x; i < sequence_length; i += blockSize) {
output_block[i] = (T)(0);
}
return;
}
// local attention token
int col_start = 0;
int col_end = sequence_length;
bool is_local_row = (global_attention[batch_index * sequence_length + row_index] == (int)0);
if (is_local_row) {
col_start = row_index - attention_window;
if (col_start < 0) {
col_start = 0;
}
col_end = row_index + attention_window + 1;
if (col_end > sequence_length) {
col_end = sequence_length;
}
}
const T* mask_block = attention_mask + sequence_length * batch_index;
int tid = threadIdx.x;
// calculate max input
float max_input = -CUDART_INF_F;
// #pragma unroll 16
for (int i = tid + col_start; i < col_end; i += blockSize) {
float x = input_block[i];
x = x * scaler + (float)mask_block[i];
if (max_input < x) {
max_input = x;
}
}
if (is_local_row) {
for (int g = tid; g < global_num; g += blockSize) {
int i = global_index[g];
if (i < col_start || i > col_end) {
float x = input_block[i];
x = x * scaler + (float)mask_block[i];
if (max_input < x) {
max_input = x;
}
}
}
}
float max_block = BlockReduce(block_reduce_temp).Reduce(max_input, cub::Max());
if (tid == 0) {
max_shared = max_block;
}
__syncthreads();
float sum_input = 0.f;
// #pragma unroll 16
for (int i = tid + col_start; i < col_end; i += blockSize) {
float x = input_block[i];
x = expf((x)*scaler + (float)mask_block[i] - max_shared);
sum_input += x;
}
if (is_local_row) {
for (int g = tid; g < global_num; g += blockSize) {
int i = global_index[g];
if (i < col_start || i > col_end) {
float x = input_block[i];
x = expf((x)*scaler + (float)mask_block[i] - max_shared);
sum_input += x;
}
}
}
float sum_block = BlockReduce(block_reduce_temp).Reduce(sum_input, cub::Sum());
if (tid == 0) {
sum_shared = sum_block;
}
__syncthreads();
float recip_sum = 1.f / sum_shared;
if (is_local_row) {
// We only need to fill in zeros for blocks that will be used in the matrix multiplication
// following the Softmax.
//
// For now zero-out only [row_index - 2*attention_window, row_index + 2*attention_window],
// we can even be more agressive and reduce the zeroing out window size since
// each row has entries in 3 blocks (3*attention_window size instead of 4*attention_window)
int zero_start = row_index - 2 * attention_window;
if (zero_start < 0) {
zero_start = 0;
}
int zero_end = row_index + 2 * attention_window;
if (zero_end > sequence_length) {
zero_end = sequence_length;
}
for (int i = tid + zero_start; i < zero_end; i += blockSize) {
output_block[i] = (T)(0.);
}
for (int g = tid; g < global_num; g += blockSize) {
int i = global_index[g];
float x = input_block[i];
x = expf((x)*scaler + (float)mask_block[i] - max_shared);
output_block[i] = (T)(recip_sum * x);
}
}
// #pragma unroll 16
for (int i = tid + col_start; i < col_end; i += blockSize) {
float x = input_block[i];
x = expf((x)*scaler + (float)mask_block[i] - max_shared);
output_block[i] = (T)(recip_sum * x);
}
}
// Launch the softmax kernel for non compact memory.
bool launchSoftmaxFastKernel(
cudaStream_t stream,
cublasHandle_t cublas,
void* workspace, // softmax space
const void* q, // transposed Q with shape (B, N, S, H)
const void* k, // transposed K with shape (B, N, S, H)
const void* v, // transposed V with shape (B, N, S, H)
const void* attention_mask, // attention mask with shape (B, S), with value 0.0 not masked, and -10000.0 masked.
const void* global_q, // Q for global tokens with shape (B, N, S, H)
const void* global_k, // K for global tokens with shape (B, N, S, H)
const void* global_v, // V for global tokens with shape (B, N, S, H)
const int* global_attention, // global attention with shape (B, S), with value 0 for local attention and 1 for global attention.
const int* global_index, // Global index with shape (B, S)
const int* batch_global_num, // Number of global tokens per batch with shape (B, 1)
void* pinned_buffer, // Pinned memory in CPU. Number of global tokens per batch with shape (B, 1)
void* output, // output with shape (B, N, S, H)
float scaler, // scalar
int batch_size, // batch size
int sequence_length, // sequence length
int num_heads, // number of heads
int head_size, // hidden size per head
int attention_window, // one sided windows size
size_t element_size) { // size of element: 2 for half, and 4 for float
bool is_fp16 = (element_size == 2);
void* scratch1 = reinterpret_cast<char*>(workspace);
void* scratch2 = reinterpret_cast<char*>(scratch1) + GetAttentionScratchSize(element_size, batch_size, num_heads, sequence_length, sequence_length);
// setup shared parameters for two strided batched matrix multiplies
cudaDataType_t Atype;
cudaDataType_t Btype;
cudaDataType_t Ctype;
cudaDataType_t resultType;
cublasGemmAlgo_t algo = CUBLAS_GEMM_DEFAULT;
__half one_fp16, zero_fp16;
float one_fp32, zero_fp32;
void *alpha, *beta_0, *beta_1;
if (is_fp16) {
one_fp16 = __float2half(1.f);
zero_fp16 = __float2half(0.f);
alpha = static_cast<void*>(&one_fp16);
beta_0 = static_cast<void*>(&zero_fp16);
beta_1 = static_cast<void*>(&one_fp16);
Atype = CUDA_R_16F;
Btype = CUDA_R_16F;
Ctype = CUDA_R_16F;
resultType = CUDA_R_16F;
algo = CUBLAS_GEMM_DEFAULT_TENSOR_OP;
} else {
one_fp32 = 1.f;
zero_fp32 = 0.f;
alpha = static_cast<void*>(&one_fp32);
beta_0 = static_cast<void*>(&zero_fp32);
beta_1 = static_cast<void*>(&one_fp32);
Atype = CUDA_R_32F;
Btype = CUDA_R_32F;
Ctype = CUDA_R_32F;
resultType = CUDA_R_32F;
}
// Strided batch matrix multiply
// qk = q * k^T
// Shapes: q and k = B x N x S x H, qk = B x N x S x S
// Convert col-major to row-major by swapping q and k in Gemm
// Local attention part
// S x S is calculated using sliding block WxW (W is one sided window size) like the following:
// [W][W]
// [W][W][W]
// [W][W][W]
// [W][W]
// The first and last rows have 2 blocks, and the remaining has 3 blocks per row.
// The calculation are splited into 3 parts. Firstly, fill the middle rows, then the first row and finally the last row.
// The results are stored in scratch1.
int w = attention_window;
int x_offset = num_heads * sequence_length * head_size;
int y_offset = num_heads * sequence_length * sequence_length;
int last_block = (sequence_length / w) - 1;
int strideA = sequence_length * head_size;
int strideB = sequence_length * head_size;
int strideC = sequence_length * sequence_length;
// When S == 2W, there is no middle rows of blocks:
// [W][W]
// [W][W]
// We can use normal matrix multiplication in this case.
if (sequence_length == 2 * w) {
CHECK(cublasGemmStridedBatchedEx(cublas,
CUBLAS_OP_T,
CUBLAS_OP_N,
sequence_length,
sequence_length,
head_size,
alpha,
k,
Atype,
head_size,
sequence_length * head_size,
q,
Btype,
head_size,
sequence_length * head_size,
beta_0,
scratch1,
Ctype,
sequence_length,
sequence_length * sequence_length,
batch_size * num_heads,
resultType,
algo));
} else { // sequence_length > 2 * w
for (int i = 0; i < batch_size; ++i) {
for (int j = 0; j < num_heads; ++j) {
void* q_head = (char*)q + (i * x_offset + j * sequence_length * head_size + w * head_size) * element_size;
void* k_head = (char*)k + (i * x_offset + j * sequence_length * head_size) * element_size;
void* qk_head = (char*)scratch1 + (i * y_offset + j * sequence_length * sequence_length + w * sequence_length) * element_size;
int count = (sequence_length - 2 * w) / w;
CHECK(cublasGemmStridedBatchedEx(cublas,
CUBLAS_OP_T,
CUBLAS_OP_N,
3 * w, // m
w, // n
head_size, // k
alpha, // alpha
k_head, // A
Atype, // A type
head_size, // lda
w * head_size, // strideA
q_head, // B
Btype, // B type
head_size, // ldb
w * head_size, // strideB
beta_0, // beta
qk_head, // C
Ctype, // C type
sequence_length, // ldc
sequence_length * w + w, // strideC
count, // batch count
resultType,
algo));
}
}
CHECK(cublasGemmStridedBatchedEx(cublas,
CUBLAS_OP_T,
CUBLAS_OP_N,
2 * w, // m
w, // n
head_size, // k
alpha, // alpha
k, // A
Atype, // A type
head_size, // lda
strideA, // strideA
q, // B
Btype, // B type
head_size, // ldb
strideB, // strideB
beta_0, // beta
scratch1, // C
Ctype, // C type
sequence_length, // ldc
strideC, // strideC
batch_size * num_heads, // batch count
resultType,
algo));
void* q_head = (char*)q + (last_block * w * head_size) * element_size;
void* k_head = (char*)k + ((last_block - 1) * w * head_size) * element_size;
void* qk_head = (char*)scratch1 + (last_block * w * sequence_length + (last_block - 1) * w) * element_size;
CHECK(cublasGemmStridedBatchedEx(cublas,
CUBLAS_OP_T,
CUBLAS_OP_N,
2 * w,
w,
head_size,
alpha,
k_head,
Atype,
head_size,
strideA,
q_head,
Btype,
head_size,
strideB,
beta_0,
qk_head,
Ctype,
sequence_length,
strideC,
batch_size * num_heads,
resultType,
algo));
}
const int* batch_global_count = reinterpret_cast<const int*>(pinned_buffer);
// Global attention part
for (int i = 0; i < batch_size; ++i) {
if (batch_global_count[i] > 0) {
void* q_batch = (char*)q + (i * x_offset) * element_size;
void* k_batch = (char*)k + (i * x_offset) * element_size;
void* qk_batch = (char*)scratch1 + (i * y_offset) * element_size;
// Local tokens attending global tokens
CHECK(cublasGemmStridedBatchedEx(cublas,
CUBLAS_OP_T,
CUBLAS_OP_N,
batch_global_count[i],
sequence_length,
head_size,
alpha,
k_batch,
Atype,
head_size,
strideA,
q_batch,
Btype,
head_size,
strideB,
beta_0,
qk_batch,
Ctype,
sequence_length,
strideC,
num_heads,
resultType,
algo));
void* global_q_batch = (char*)global_q + (i * num_heads * sequence_length * head_size) * element_size;
void* global_k_batch = (char*)global_k + (i * x_offset) * element_size;
int strideB_global = sequence_length * head_size;
// Global tokens attending everything
// This GEMMs need to be last to make sure all global token entries are re-written.
CHECK(cublasGemmStridedBatchedEx(cublas,
CUBLAS_OP_T,
CUBLAS_OP_N,
sequence_length,
batch_global_count[i],
head_size,
alpha,
global_k_batch,
Atype,
head_size,
strideA,
global_q_batch,
Btype,
head_size,
strideB_global,
beta_0,
qk_batch,
Ctype,
sequence_length,
strideC,
num_heads,
resultType,
algo));
}
}
int dim0 = sequence_length * num_heads;
int dim1 = sequence_length;
void* softmax_out = scratch2;
const int blockSize = 64;
const int gridSize = batch_size * num_heads * sequence_length;
if (is_fp16) {
LongformerSoftmaxFastKernel<__half, blockSize><<<gridSize, blockSize, 0, stream>>>(
global_attention,
global_index,
batch_global_num,
static_cast<const __half*>(scratch1),
static_cast<const __half*>(attention_mask),
static_cast<__half*>(softmax_out), scaler, dim0, dim1, attention_window);
} else {
LongformerSoftmaxFastKernel<float, blockSize><<<gridSize, blockSize, 0, stream>>>(
global_attention,
global_index,
batch_global_num,
static_cast<const float*>(scratch1),
static_cast<const float*>(attention_mask),
static_cast<float*>(softmax_out), scaler, dim0, dim1, attention_window);
}
// Run the matrix multiply: output = softmax_out * v
// softmax_out: B x N x S x S
// v: B x N x S x H
// attn_out: B x N x S x H
// Calculation uses full Gemm (S == 2W) or sliding blocks (S > 2W) in a way similar to local attention part.
if (sequence_length == 2 * w) {
// convert col-major to row-major by swapping softmax_out and v
CHECK(cublasGemmStridedBatchedEx(cublas,
CUBLAS_OP_N,
CUBLAS_OP_N,
head_size,
sequence_length,
sequence_length,
alpha,
v,
Atype,
head_size,
sequence_length * head_size,
softmax_out,
Btype,
sequence_length,
sequence_length * sequence_length,
beta_0,
output,
Ctype,
head_size,
sequence_length * head_size,
batch_size * num_heads,
resultType,
algo));
} else { // sequence_length > 2 * w
for (int i = 0; i < batch_size; ++i) {
for (int j = 0; j < num_heads; ++j) {
void* v_head = (char*)v + (i * x_offset + j * head_size * sequence_length) * element_size;
void* prob_head = (char*)softmax_out + (i * y_offset + j * sequence_length * sequence_length + w * sequence_length) * element_size;
void* out_head = (char*)output + (i * x_offset + j * head_size * sequence_length + w * head_size) * element_size;
int count = (sequence_length - 2 * w) / w;
CHECK(cublasGemmStridedBatchedEx(cublas,
CUBLAS_OP_N,
CUBLAS_OP_N,
head_size,
w,
3 * w,
alpha,
v_head,
Atype,
head_size,
w * head_size,
prob_head,
Btype,
sequence_length,
sequence_length * w + w,
beta_0,
out_head,
Ctype,
head_size,
w * head_size,
count,
resultType,
algo));
}
}
CHECK(cublasGemmStridedBatchedEx(cublas,
CUBLAS_OP_N,
CUBLAS_OP_N,
head_size,
w,
2 * w,
alpha,
v,
Atype,
head_size,
sequence_length * head_size,
softmax_out,
Btype,
sequence_length,
sequence_length * sequence_length,
beta_0,
output,
Ctype,
head_size,
sequence_length * head_size,
batch_size * num_heads,
resultType,
algo));
void* v_head = (char*)v + (last_block - 1) * w * head_size * element_size;
void* prob_head = (char*)softmax_out + (sequence_length * last_block * w + (last_block - 1) * w) * element_size;
void* out_head = (char*)output + last_block * w * head_size * element_size;
CHECK(cublasGemmStridedBatchedEx(cublas,
CUBLAS_OP_N,
CUBLAS_OP_N,
head_size,
w,
2 * w,
alpha,
v_head,
Atype,
head_size,
sequence_length * head_size,
prob_head,
Btype,
sequence_length,
sequence_length * sequence_length,
beta_0,
out_head,
Ctype,
head_size,
sequence_length * head_size,
batch_size * num_heads,
resultType,
algo));
}
for (int i = 0; i < batch_size; ++i) {
if (batch_global_count[i] > 0) {
int glob_longdim_mm = (last_block - 1) * w;
void* v_head = (char*)v + (i * x_offset) * element_size;
void* prob_head = (char*)softmax_out + (i * y_offset + 2 * w * sequence_length) * element_size;
void* out_head = (char*)output + (i * x_offset + 2 * w * head_size) * element_size;
CHECK(cublasGemmStridedBatchedEx(cublas,
CUBLAS_OP_N,
CUBLAS_OP_N,
head_size,
glob_longdim_mm,
batch_global_count[i],
alpha,
v_head,
Atype,
head_size,
sequence_length * head_size,
prob_head,
Btype,
sequence_length,
sequence_length * sequence_length,
beta_1,
out_head,
Ctype,
head_size,
sequence_length * head_size,
num_heads,
resultType,
algo));
// Global tokens
v_head = (char*)global_v + (i * x_offset) * element_size;
prob_head = (char*)softmax_out + (i * y_offset) * element_size;
out_head = (char*)output + (i * x_offset) * element_size;
CHECK(cublasGemmStridedBatchedEx(cublas,
CUBLAS_OP_N,
CUBLAS_OP_N,
head_size,
batch_global_count[i],
sequence_length, // Re-write entries completely
alpha,
v_head,
Atype,
head_size,
sequence_length * head_size,
prob_head,
Btype,
sequence_length,
sequence_length * sequence_length,
beta_0, // Use beta=0 to overwrite
out_head, // Here assumes global tokens are at the beginning of sequence.
Ctype,
head_size,
sequence_length * head_size,
num_heads,
resultType,
algo));
}
}
return true;
}
} // namespace cuda
} // namespace contrib
} // namespace onnxruntime

51
onnxruntime/contrib_ops/cuda/bert/longformer_attention_softmax.h Normal file
View file

@ -0,0 +1,51 @@
/*
Copyright (c) NVIDIA Corporation and Microsoft Corporation
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 is fast cuda kernels for longformer attention softmax.
// It uses two temporary matrix of BxNxSxS, and consumes more memory when sequence length is large.
namespace onnxruntime {
namespace contrib {
namespace cuda {
// Launch the softmax kernel for non compact memory.
bool launchSoftmaxFastKernel(
cudaStream_t stream,
cublasHandle_t cublas,
void* workspace, // softmax space
const void* q, // transposed Q with shape (B, N, S, H)
const void* k, // transposed K with shape (B, N, S, H)
const void* v, // transposed V with shape (B, N, S, H)
const void* attention_mask, // attention mask with shape (B, S), with value 0.0 not masked, and -10000.0 masked.
const void* global_q, // Q for global tokens with shape (B, N, S, H)
const void* global_k, // K for global tokens with shape (B, N, S, H)
const void* global_v, // V for global tokens with shape (B, N, S, H)
const int* global_attention, // global attention with shape (B, S), with value 0 for local attention and 1 for global attention.
const int* global_index, // Global index with shape (B, S)
const int* batch_global_num, // Number of global tokens per batch with shape (B, 1)
void* pinned_buffer, // Pinned memory in CPU. Number of global tokens per batch with shape (B, 1)
void* output, // output with shape (B, N, S, H)
float scaler, // scalar
int batch_size, // batch size
int sequence_length, // sequence length
int num_heads, // number of heads
int head_size, // hidden size per head
int attention_window, // one sided windows size
size_t element_size);
} // namespace cuda
} // namespace contrib
} // namespace onnxruntime

11
onnxruntime/core/optimizer/constant_folding.cc
View file

@ -182,9 +182,16 @@ Status ConstantFolding::ApplyImpl(Graph& graph, bool& modified, int graph_level,
if (converted_to_constant) {
// Remove single-output node chain for inputs of the node
for (auto p_ip_node = node->InputNodesBegin(); p_ip_node != node->InputNodesEnd(); ++p_ip_node) {
graph_utils::RemoveNodesWithOneOutputBottomUp(graph, *p_ip_node);
auto p_ip_node = node->InputNodesBegin();
const auto p_ip_node_end = node->InputNodesEnd();
while (p_ip_node != p_ip_node_end) {
const auto& input_node = *p_ip_node;
// Update the node iterator before removing the corresponding node because removing
// the node will invalidate the node iterator
++p_ip_node;
graph_utils::RemoveNodesWithOneOutputBottomUp(graph, input_node);
}
// Remove the output edges of the constant node and then remove the node itself.
graph_utils::RemoveNodeOutputEdges(graph, *node);
graph.RemoveNode(node->Index());

42
onnxruntime/core/providers/coreml/builders/helper.cc
View file

@ -3,6 +3,11 @@
#include <vector>
#ifdef __APPLE__
#include <sys/utsname.h>
#include <TargetConditionals.h>
#endif
#include "helper.h"
#include <core/graph/graph_viewer.h>
@ -122,5 +127,42 @@ std::vector<std::vector<NodeIndex>> GetSupportedNodes(const GraphViewer& graph_v
return supported_node_vecs;
}
bool HasNeuralEngine(const logging::Logger& logger) {
bool has_neural_engine = false;
#ifdef __APPLE__
struct utsname system_info;
uname(&system_info);
LOGS(logger, VERBOSE) << "Current Apple hardware info: " << system_info.machine;
#if TARGET_OS_IPHONE
// utsname.machine has device identifier. For example, identifier for iPhone Xs is "iPhone11,2".
// Since Neural Engine is only available for use on A12 and later, major device version in the
// identifier is checked for these models:
// A12: iPhone XS (11,2), iPad Mini - 5th Gen (11,1)
// A12X: iPad Pro - 3rd Gen (8,1)
// For more information, see https://www.theiphonewiki.com/wiki/Models
size_t str_len = strlen(system_info.machine);
if (str_len > 4 && strncmp("iPad", system_info.machine, 4) == 0) {
const int major_version = atoi(system_info.machine + 4);
has_neural_engine = major_version >= 8; // There are no device between iPad 8 and 11.
} else if (str_len > 6 && strncmp("iPhone", system_info.machine, 6) == 0) {
const int major_version = atoi(system_info.machine + 6);
has_neural_engine = major_version >= 11;
}
#elif TARGET_OS_OSX && TARGET_CPU_ARM64
// Only Mac with arm64 CPU (Apple Silicon) has ANE.
has_neural_engine = true;
#endif // #if TARGET_OS_IPHONE
#else
// In this case, we are running the EP on non-apple platform, which means we are running the model
// conversion with CoreML EP enabled, for this we always assume the target system has Neural Engine
LOGS(logger, VERBOSE) << "HasNeuralEngine running on non-Apple hardware for model conversion only";
has_neural_engine = true;
#endif // #ifdef __APPLE__
return has_neural_engine;
}
} // namespace coreml
} // namespace onnxruntime

4
onnxruntime/core/providers/coreml/builders/helper.h
View file

@ -28,5 +28,9 @@ bool IsInputSupported(const NodeArg& node_arg, const std::string& parent_name, c
std::vector<std::vector<NodeIndex>> GetSupportedNodes(const GraphViewer& graph_viewer,
const logging::Logger& logger);
// CoreML is more efficient running using Apple Neural Engine
// This is to detect if the current system has Apple Neural Engine
bool HasNeuralEngine(const logging::Logger& logger);
} // namespace coreml
} // namespace onnxruntime

7
onnxruntime/core/providers/coreml/coreml_execution_provider.cc
View file

@ -67,6 +67,13 @@ CoreMLExecutionProvider::GetCapability(const onnxruntime::GraphViewer& graph_vie
*/
const auto& logger = *GetLogger();
bool has_neural_engine = coreml::HasNeuralEngine(logger);
if ((coreml_flags_ & COREML_FLAG_ONLY_ENABLE_DEVICE_WITH_ANE) && !has_neural_engine) {
LOGS(logger, VERBOSE) << "The current system does not have Apple Neural Engine";
return result;
}
const auto node_groups = coreml::GetSupportedNodes(graph_viewer, logger);
if (node_groups.empty()) {

13
onnxruntime/core/providers/cpu/math/einsum_utils/einsum_auxiliary_ops.cc
View file

@ -456,5 +456,18 @@ template std::unique_ptr<Tensor> ReduceSum<int64_t>(
const std::vector<int64_t>& reduce_axes, AllocatorPtr allocator,
concurrency::ThreadPool* tp, void* einsum_cuda_assets, const DeviceHelpers::ReduceSum<int64_t>& reduce_sum_func);
// MLFloat16
template std::unique_ptr<Tensor> MatMul<MLFloat16>(
const Tensor& input_1, const std::vector<int64_t>& input_shape_1_override,
const Tensor& input_2, const std::vector<int64_t>& input_shape_2_override,
AllocatorPtr allocator, concurrency::ThreadPool* tp, void* einsum_cuda_assets,
const DeviceHelpers::MatMul<MLFloat16>& device_matmul_func);
template std::unique_ptr<Tensor> ReduceSum<MLFloat16>(
const Tensor& input, const std::vector<int64_t>& input_shape_override,
const std::vector<int64_t>& reduce_axes, AllocatorPtr allocator,
concurrency::ThreadPool* tp, void* einsum_cuda_assets,
const DeviceHelpers::ReduceSum<MLFloat16>& device_reduce_sum_func);
} // namespace EinsumOp
} // namespace onnxruntime

1
onnxruntime/core/providers/cpu/math/einsum_utils/einsum_typed_compute_processor.cc
View file

@ -367,5 +367,6 @@ template class EinsumTypedComputeProcessor<float>;
template class EinsumTypedComputeProcessor<int32_t>;
template class EinsumTypedComputeProcessor<double>;
template class EinsumTypedComputeProcessor<int64_t>;
template class EinsumTypedComputeProcessor<MLFloat16>;
} // namespace onnxruntime

2
onnxruntime/core/providers/cuda/cuda_allocator.h
View file

@ -28,7 +28,7 @@ class CUDAExternalAllocator : public CUDAAllocator {
typedef void (*ExternalFree)(void* p);
public:
CUDAExternalAllocator(OrtDevice::DeviceId device_id, const char* name, void* alloc, void* free)
CUDAExternalAllocator(OrtDevice::DeviceId device_id, const char* name, const void* alloc, const void* free)
: CUDAAllocator(device_id, name) {
alloc_ = reinterpret_cast<ExternalAlloc>(alloc);
free_ = reinterpret_cast<ExternalFree>(free);

10
onnxruntime/core/providers/cuda/cuda_execution_provider.cc
View file

@ -127,11 +127,19 @@ CUDAExecutionProvider::CUDAExecutionProvider(const CUDAExecutionProviderInfo& in
// must wait GPU idle, otherwise cudaGetDeviceProperties might fail
CUDA_CALL_THROW(cudaDeviceSynchronize());
CUDA_CALL_THROW(cudaGetDeviceProperties(&device_prop_, info_.device_id));
// This scenario is not supported.
ORT_ENFORCE(!(info.has_user_compute_stream && info.external_allocator_info.UseExternalAllocator()));
if (info.has_user_compute_stream) {
external_stream_ = true;
stream_ = static_cast<cudaStream_t>(info.user_compute_stream);
} else {
CUDA_CALL_THROW(cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking));
if (info.external_allocator_info.UseExternalAllocator()) {
stream_ = nullptr;
} else {
CUDA_CALL_THROW(cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking));
}
}
size_t free = 0;

13
onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc
View file

@ -36,7 +36,8 @@ const EnumNameMapping<ArenaExtendStrategy> arena_extend_strategy_mapping{
CUDAExecutionProviderInfo CUDAExecutionProviderInfo::FromProviderOptions(const ProviderOptions& options) {
CUDAExecutionProviderInfo info{};
void* alloc = nullptr;
void* free = nullptr;
ORT_THROW_IF_ERROR(
ProviderOptionsParser{}
.AddValueParser(
@ -55,18 +56,18 @@ CUDAExecutionProviderInfo CUDAExecutionProviderInfo::FromProviderOptions(const P
})
.AddValueParser(
cuda::provider_option_names::kcudaExternalAlloc,
[&info](const std::string& value_str) -> Status {
[&alloc](const std::string& value_str) -> Status {
size_t address;
ORT_RETURN_IF_ERROR(ParseStringWithClassicLocale(value_str, address));
info.external_allocator_info.alloc = reinterpret_cast<void*>(address);
alloc = reinterpret_cast<void*>(address);
return Status::OK();
})
.AddValueParser(
cuda::provider_option_names::kcudaExternalFree,
[&info](const std::string& value_str) -> Status {
[&free](const std::string& value_str) -> Status {
size_t address;
ORT_RETURN_IF_ERROR(ParseStringWithClassicLocale(value_str, address));
info.external_allocator_info.free = reinterpret_cast<void*>(address);
free = reinterpret_cast<void*>(address);
return Status::OK();
})
.AddAssignmentToReference(cuda::provider_option_names::kMemLimit, info.cuda_mem_limit)
@ -79,6 +80,8 @@ CUDAExecutionProviderInfo CUDAExecutionProviderInfo::FromProviderOptions(const P
.AddAssignmentToReference(cuda::provider_option_names::kDoCopyInDefaultStream, info.do_copy_in_default_stream)
.Parse(options));
CUDAExecutionProviderExternalAllocatorInfo alloc_info{alloc, free};
info.external_allocator_info = alloc_info;
return info;
}

11
onnxruntime/core/providers/cuda/cuda_execution_provider_info.h
View file

@ -13,15 +13,20 @@
namespace onnxruntime {
// Information needed to construct CUDA execution providers.
struct CUDAExecutionProviderExternalAllocatorInfo {
void* alloc{nullptr};
void* free{nullptr};
const void* alloc{nullptr};
const void* free{nullptr};
CUDAExecutionProviderExternalAllocatorInfo() {
alloc = nullptr;
free = nullptr;
}
bool UseExternalAllocator() {
CUDAExecutionProviderExternalAllocatorInfo(void* a, void* f) {
alloc = a;
free = f;
}
bool UseExternalAllocator() const {
return (alloc != nullptr) && (free != nullptr);
}
};

13
onnxruntime/core/providers/cuda/math/einsum.cc
View file

@ -15,7 +15,8 @@ ONNX_OPERATOR_KERNEL_EX(
KernelDefBuilder().TypeConstraint("T",
std::vector<MLDataType>{
DataTypeImpl::GetTensorType<float>(),
DataTypeImpl::GetTensorType<double>()}),
DataTypeImpl::GetTensorType<double>(),
DataTypeImpl::GetTensorType<MLFloat16>()}),
Einsum);
Status Einsum::Compute(OpKernelContext* context) const {
@ -59,6 +60,16 @@ Status Einsum::DeviceCompute(OpKernelContext* context, const std::vector<const T
EinsumOp::DeviceHelpers::CudaDeviceHelpers::ReduceSum<double>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::DataCopy);
return einsum_compute_processor.Run();
} else if (inputs[0]->IsDataType<MLFloat16>()) {
auto einsum_compute_processor = EinsumTypedComputeProcessor<MLFloat16>(context, allocator, tp,
einsum_compute_preprocessor,
&einsum_cuda_assets);
einsum_compute_processor.SetDeviceHelpers(EinsumOp::DeviceHelpers::CudaDeviceHelpers::Transpose,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::MatMul<MLFloat16>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::ReduceSum<MLFloat16>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::DataCopy);
return einsum_compute_processor.Run();
}
return ORT_MAKE_STATUS(ONNXRUNTIME, NOT_IMPLEMENTED,

13
onnxruntime/core/providers/cuda/math/einsum_utils/einsum_auxiliary_ops.cc
View file

@ -165,6 +165,19 @@ template Tensor DeviceHelpers::CudaDeviceHelpers::ReduceSum<double>(
const TensorShape* input_shape_override,
concurrency::ThreadPool* tp, void* einsum_cuda_assets);
// MLFloat16
template Status DeviceHelpers::CudaDeviceHelpers::MatMul<MLFloat16>(
const MLFloat16* input_1_data, const MLFloat16* input_2_data, MLFloat16* output_data,
size_t left_stride, size_t right_stride, size_t output_stride,
size_t num_batches, size_t M, size_t K, size_t N, concurrency::ThreadPool* tp,
void* einsum_cuda_assets);
template Tensor DeviceHelpers::CudaDeviceHelpers::ReduceSum<MLFloat16>(
const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
concurrency::ThreadPool* tp, void* einsum_cuda_assets);
} // namespace EinsumOp
} // namespace onnxruntime

11
onnxruntime/core/providers/cuda/math/einsum_utils/einsum_auxiliary_ops_diagonal.cu
View file

@ -30,8 +30,8 @@ __global__ void _DiagonalKernel(
if (i == dim_1) {
// Process dim_2 as dim_2 needs to have the same dim value as dim_1
// For example: given a tensor of shape [2, 3, 3] and parsing the diagonal along axes `1` and `2`
// we need to parse elements in input[j, i, i] (j -> 0 to 1; and i -> 0 to 2)
// and place them in output[j, i] and by definition of diagonal parsing dim_1 has to be equal to
// we need to parse elements in input[j, i, i] (j -> 0 to 1; and i -> 0 to 2)
// and place them in output[j, i] and by definition of diagonal parsing dim_1 has to be equal to
// dim_2
input_idx += input_strides[dim_2] * dim;
}
@ -75,6 +75,13 @@ void DiagonalImpl(
output_size);
break;
case sizeof(int16_t):
_DiagonalKernel<half><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0, stream>>>(
reinterpret_cast<const half*>(input_data), input_rank, dim_1, dim_2,
input_strides, reinterpret_cast<half*>(output_data), output_strides,
output_size);
break;
// Should not hit this as we do not register kernel support for types that will run into this
default:
ORT_THROW("Einsum Op: Diagonal parsing unsupported");

5
onnxruntime/core/providers/cuda/nn/conv_transpose.cc
View file

@ -45,10 +45,9 @@ Status ConvTranspose<T>::DoConvTranspose(OpKernelContext* context, bool dynamic_
auto x_data = reinterpret_cast<const CudaT*>(X->template Data<T>());
auto x_dimensions = X->Shape().NumDimensions();
if (x_dimensions != 4 && x_dimensions != 3) {
// This condition is not true for test_convtranspose_3d in ONNX tests series.
if (x_dimensions < 3 || x_dimensions > 5) {
// TODO: the error message should tell which operator raises it.
return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Input X must be 3- or 4-dimensional.",
return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Input X must be 3-, 4- or 5-dimensional.",
" X: ", X->Shape().ToString().c_str());
}
const Tensor* W = context->Input<Tensor>(1);

9
onnxruntime/core/providers/cuda/reduction/reduction_ops.cc
View file

@ -610,7 +610,7 @@ Status ReduceComputeCore(CUDAExecutionProvider& cuda_ep, const Tensor& input, Pr
&zero, output_tensor, reinterpret_cast<CudaT*>(output.template MutableData<T>())));
}
}
} else {
} else {
// For ArgMax & ArgMin ops, use the indicies as the output with int64 type
// cudnnReduceTensor has issue if input and output has same size, which will happen if the axis to be reduced has dim value of 1.
// the output is zeros of the output size
@ -928,6 +928,13 @@ template Tensor ReduceCompute<double, CUDNN_REDUCE_TENSOR_NO_INDICES>(
bool keep_dims, bool calculate_log, bool calculate_sqt, bool log_sum_exp,
bool fast_reduction, const TensorShape* input_shape_override);
template Tensor ReduceCompute<MLFloat16, CUDNN_REDUCE_TENSOR_NO_INDICES>(
CUDAExecutionProvider& cuda_ep, cudnnReduceTensorOp_t cudnn_reduce_op,
AllocatorPtr allocator,
const Tensor& input, const std::vector<int64_t>& axes,
bool keep_dims, bool calculate_log, bool calculate_sqt, bool log_sum_exp,
bool fast_reduction, const TensorShape* input_shape_override);
} // namespace ReductionOps
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000

2
onnxruntime/core/providers/op_kernel_type_control_overrides.inc
View file

@ -47,7 +47,7 @@ namespace op_kernel_type_control {
// ORT_SPECIFY_OP_KERNEL_ARG_ALLOWED_TYPES(kOnnxDomain, Cast, Input, 0, float, int64_t);
// ORT_SPECIFY_OP_KERNEL_ARG_ALLOWED_TYPES(kOnnxDomain, Cast, Output, 0, float, int64_t);
// Specify allowed types globally:
// ORT_SPECIFY_OP_KERNEL_GLOBAL_ALLOWED_TYPES(float, double, int32_t)
// ORT_SPECIFY_OP_KERNEL_GLOBAL_ALLOWED_TYPES(float, double, int32_t);
// specify allowed types here

106
onnxruntime/core/providers/rocm/rocm_execution_provider.cc
View file

@ -674,7 +674,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, MLFloat16, ArgMin);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, Compress);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, Concat);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, Flatten);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, Flatten);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, Gather);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, GatherElements);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, float, Gemm);
@ -767,7 +767,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, uint8_t, QuantizeLinear);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, int8_t, DequantizeLinear);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, uint8_t, DequantizeLinear);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, CumSum);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 13, CumSum);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, int64_t_int64_t_int64_t, OneHot);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, int64_t_float_int64_t, OneHot);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, int32_t_float_int32_t, OneHot);
@ -775,7 +775,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, int32_t_MLFloat16_int32_t, OneHot);
// OpSet 12
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, Clip);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, 12, Clip);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, float, MaxPool);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, double, MaxPool);
@ -989,6 +989,19 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, MLFloat16, ReduceSumSquare);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, int64_t, GatherND);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, Dropout);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, float, Resize);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, double, Resize);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, MLFloat16, Resize);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, int32_t, Resize);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, uint8_t, Resize);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, If);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, Loop);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, Flatten);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, float, LRN);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, double, LRN);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, MLFloat16, LRN);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, Identity);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, ScatterND);
template <>
KernelCreateInfo BuildKernelCreateInfo<void>() {
@ -1252,9 +1265,9 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, 12, uint32_t, Cast)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, 12, uint64_t, Cast)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, 12, bool, Cast)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 2, 10, float, Pad)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 2, 10, double, Pad)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 2, 10, MLFloat16, Pad)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 2, 10, float, Pad)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 2, 10, double, Pad)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 2, 10, MLFloat16, Pad)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 1, 4, Reshape)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 5, 12, Reshape)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 1, 12, Shape)>,
@ -1277,26 +1290,26 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 1, 9, int32_t, Slice)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 1, 9, int64_t, Slice)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 1, 9, float, Slice)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, 10, Compress)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, 10, Flatten)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, 10, Compress)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, 10, Flatten)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 7, 9, float, Upsample)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 7, 9, double, Upsample)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 7, 9, MLFloat16, Upsample)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 7, 9, int32_t, Upsample)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 7, 9, uint8_t, Upsample)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 2, 10, Split)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, ConstantOfShape)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, int8_t, Shrink)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, int16_t, Shrink)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, int32_t, Shrink)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, int64_t, Shrink)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, uint8_t, Shrink)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, uint16_t, Shrink)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, uint32_t, Shrink)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, uint64_t, Shrink)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, float, Shrink)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, double, Shrink)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, MLFloat16, Shrink)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 2, 10, Split)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, ConstantOfShape)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, int8_t, Shrink)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, int16_t, Shrink)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, int32_t, Shrink)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, int64_t, Shrink)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, uint8_t, Shrink)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, uint16_t, Shrink)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, uint32_t, Shrink)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, uint64_t, Shrink)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, float, Shrink)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, double, Shrink)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, MLFloat16, Shrink)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 7, 8, float, Less)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 7, 8, double, Less)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 7, 8, MLFloat16, Less)>,
@ -1307,7 +1320,7 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, 12, float, Less)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, 12, double, Less)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, 12, MLFloat16, Less)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, EyeLike)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, EyeLike)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, 10, Scatter)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, MLFloat16, Where)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 9, float, Where)>,
@ -1328,7 +1341,7 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) {
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, float, AveragePool)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, double, AveragePool)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, MLFloat16, AveragePool)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 11, Dropout)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 11, Dropout)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, float, MaxPool)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, double, MaxPool)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, MLFloat16, MaxPool)>,
@ -1338,12 +1351,12 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) {
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, MLFloat16, Resize)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, int32_t, Resize)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, uint8_t, Resize)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, ReverseSequence)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, ReverseSequence)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, float, RoiAlign)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, double, RoiAlign)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, int32_t, Slice)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, int64_t, Slice)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, float, Slice)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, int32_t, Slice)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, int64_t, Slice)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, 10, float, Slice)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, float, ThresholdedRelu)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, double, ThresholdedRelu)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 10, MLFloat16, ThresholdedRelu)>,
@ -1361,9 +1374,9 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, float, ArgMin)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, double, ArgMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, MLFloat16, ArgMin)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, Compress)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, Compress)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, Concat)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, Flatten)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, Flatten)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, Gather)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, GatherElements)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, MLFloat16, Gemm)>,
@ -1372,7 +1385,7 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) {
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, If)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, Loop)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, NonMaxSuppression)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, Range)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, Range)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, float, ReduceL1)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, double, ReduceL1)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, MLFloat16, ReduceL1)>,
@ -1411,7 +1424,7 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) {
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, double, ReduceSumSquare)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, MLFloat16, ReduceSumSquare)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, Scan)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, ScatterElements)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, ScatterElements)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, int32_t, Slice)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, int64_t, Slice)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, float, Slice)>,
@ -1442,17 +1455,17 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) {
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, MLFloat16, Resize)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, int32_t, Resize)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, uint8_t, Resize)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 11, Clip)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, float, Pad)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, double, Pad)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, MLFloat16, Pad)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 11, Clip)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, float, Pad)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, double, Pad)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, MLFloat16, Pad)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, bool, Equal)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, int32_t, Equal)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, int64_t, Equal)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, float, Round)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, double, Round)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, MLFloat16, Round)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, CumSum)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, float, Round)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, double, Round)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, MLFloat16, Round)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 13, CumSum)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, int64_t_int64_t_int64_t, OneHot)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, int64_t_float_int64_t, OneHot)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, int32_t_float_int32_t, OneHot)>,
@ -1460,7 +1473,7 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) {
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, int32_t_MLFloat16_int32_t, OneHot)>,
// OpSet 12
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, Clip)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, 12, Clip)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, float, MaxPool)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, double, MaxPool)>,
@ -1674,6 +1687,19 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, MLFloat16, ReduceSumSquare)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, int64_t, GatherND)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, Dropout)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, float, Resize)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, double, Resize)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, MLFloat16, Resize)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, int32_t, Resize)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, uint8_t, Resize)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, If)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, Loop)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, Flatten)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, float, LRN)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, double, LRN)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, MLFloat16, LRN)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, Identity)>,
// BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, ScatterND)>,
};
for (auto& function_table_entry : function_table) {

3
onnxruntime/core/session/onnxruntime_c_api.cc
View file

@ -2104,12 +2104,11 @@ static constexpr OrtApi ort_api_1_to_7 = {
&OrtApis::ReleaseArenaCfg,
// End of Version 6 - DO NOT MODIFY ABOVE (see above text for more information)
// Version 7 - In development, feel free to add/remove/rearrange here
&OrtApis::ModelMetadataGetGraphDescription,
&OrtApis::SessionOptionsAppendExecutionProvider_TensorRT,
&OrtApis::SetCurrentGpuDeviceId,
&OrtApis::GetCurrentGpuDeviceId,
// End of Version 7 - DO NOT MODIFY ABOVE (see above text for more information)
};
// Assert to do a limited check to ensure Version 1 of OrtApi never changes (will detect an addition or deletion but not if they cancel out each other)

64
onnxruntime/python/_pybind_state.py
View file

@ -2,25 +2,53 @@
# Copyright (c) Microsoft Corporation. All rights reserved.
# Licensed under the MIT License.
# --------------------------------------------------------------------------
"""
Ensure that dependencies are available and then load the extension module.
"""
import os
import platform
import warnings
import onnxruntime.capi._ld_preload # noqa: F401
import sys
try:
from onnxruntime.capi.onnxruntime_pybind11_state import * # noqa
except ImportError as e:
warnings.warn("Cannot load onnxruntime.capi. Error: '{0}'.".format(str(e)))
from . import _ld_preload # noqa: F401
# If on Windows, check if this import error is caused by the user not installing the 2019 VC Runtime
# The VC Redist installer usually puts the VC Runtime dlls in the System32 folder
# This may not always paint the true picture as anyone building from source using VS 2017 might hit this error
# because the machine might be missing the 2019 VC Runtime but it is not actually needed in that case and the
# import error might actually be due to some other reason.
# TODO: Add a guard against False Positive error message
# As a proxy for checking if the 2019 VC Runtime is installed,
# we look for a specific dll only shipped with the 2019 VC Runtime
if platform.system().lower() == 'windows' and not os.path.isfile('c:\\Windows\\System32\\vcruntime140_1.dll'):
warnings.warn("Unless you have built the wheel using VS 2017, "
"please install the 2019 Visual C++ runtime and then try again")
if platform.system() == "Windows":
from . import version_info
if version_info.use_cuda:
cuda_version_major, cuda_version_minor = version_info.cuda_version.split(".")
if int(cuda_version_major) < 11:
# Prior to CUDA 11 both major and minor version at build time/runtime have to match.
cuda_env_variable = f"CUDA_PATH_V{cuda_version_major}_{cuda_version_minor}"
if cuda_env_variable not in os.environ:
raise ImportError(f"CUDA Toolkit {version_info.cuda_version} not installed on the machine.")
else:
# With CUDA 11 and newer only the major version at build time/runtime has to match.
# Use the most recent minor version available.
cuda_env_variable = None
for i in range(9, -1, -1):
if f"CUDA_PATH_V{cuda_version_major}_{i}" in os.environ:
cuda_env_variable = f"CUDA_PATH_V{cuda_version_major}_{i}"
break
if not cuda_env_variable:
raise ImportError(f"CUDA Toolkit {cuda_version_major}.x not installed on the machine.")
cuda_bin_dir = os.path.join(os.environ[cuda_env_variable], "bin")
if not os.path.isfile(os.path.join(cuda_bin_dir, f"cudnn64_{version_info.cudnn_version}.dll")):
raise ImportError(f"cuDNN {version_info.cudnn_version} not installed in {cuda_bin_dir}.")
if sys.version_info >= (3, 8):
# Python 3.8 (and later) doesn't search system PATH when loading DLLs, so the CUDA location needs to be
# specified explicitly using the new API introduced in Python 3.8.
os.add_dll_directory(cuda_bin_dir)
else:
# Python 3.7 (and earlier) searches directories listed in PATH variable.
# Make sure that the target CUDA version is at the beginning (important if multiple CUDA versions are
# installed on the machine.)
os.environ["PATH"] += cuda_bin_dir + os.pathsep + os.environ["PATH"]
if version_info.vs2019 and platform.architecture()[0] == "64bit":
if not os.path.isfile("C:\\Windows\\System32\\vcruntime140_1.dll"):
raise ImportError(
"Microsoft Visual C++ Redistributable for Visual Studio 2019 not installed on the machine.")
from .onnxruntime_pybind11_state import * # noqa

5
onnxruntime/python/tools/quantization/E2E_example_model/image_classification/cpu/run.py
View file

@ -12,7 +12,7 @@ from PIL import Image
import onnx
import onnxruntime
from onnx import helper, TensorProto, numpy_helper
from onnxruntime.quantization import quantize_static, CalibrationDataReader, QuantFormat
from onnxruntime.quantization import quantize_static, CalibrationDataReader, QuantFormat, QuantType
class ResNet50DataReader(CalibrationDataReader):
@ -107,7 +107,8 @@ def main():
output_model_path,
dr,
quant_format=args.quant_format,
per_channel=args.per_channel)
per_channel=args.per_channel,
weight_type=QuantType.QInt8)
print('Calibrated and quantized model saved.')
print('benchmarking fp32 model...')

17
onnxruntime/python/tools/quantization/E2E_example_model/image_classification/trt/resnet50/e2e_tensorrt_resnet_example.py
View file

@ -7,7 +7,7 @@ import logging
from PIL import Image
import onnx
import onnxruntime
from onnxruntime.quantization import CalibrationDataReader, create_calibrater, write_calibration_table
from onnxruntime.quantization import CalibrationDataReader, create_calibrator, write_calibration_table
class ImageNetDataReader(CalibrationDataReader):
@ -258,9 +258,13 @@ class ImageClassificationEvaluator:
def convert_model_batch_to_dynamic(model_path):
model = onnx.load(model_path)
input = model.graph.input
input_name = input[0].name
shape = input[0].type.tensor_type.shape
initializers = [node.name for node in model.graph.initializer]
inputs = []
for node in model.graph.input:
if node.name not in initializers:
inputs.append(node)
input_name = inputs[0].name
shape = inputs[0].type.tensor_type.shape
dim = shape.dim
if not dim[0].dim_param:
dim[0].dim_param = 'N'
@ -329,6 +333,8 @@ if __name__ == '__main__':
# Generate INT8 calibration table
if calibration_table_generation_enable:
calibrator = create_calibrator(new_model_path, [], augmented_model_path=augmented_model_path)
calibrator.set_execution_providers(["CUDAExecutionProvider"])
data_reader = ImageNetDataReader(ilsvrc2012_dataset_path,
start_index=0,
end_index=calibration_dataset_size,
@ -336,9 +342,6 @@ if __name__ == '__main__':
batch_size=batch_size,
model_path=augmented_model_path,
input_name=input_name)
# For TensorRT calibration, augment all FP32 tensors (empty op_types), disable ORT graph optimization and skip quantization parameter calculation
calibrator = create_calibrater(new_model_path)
calibrator.set_execution_providers(["CUDAExecutionProvider"])
calibrator.collect_data(data_reader)
write_calibration_table(calibrator.compute_range())

17
onnxruntime/python/tools/quantization/E2E_example_model/object_detection/trt/yolov3/data_reader.py
View file

@ -1,5 +1,5 @@
from onnxruntime.quantization import CalibrationDataReader
from preprocessing import yolov3_preprocess_func, yolov3_variant_preprocess_func
from preprocessing import yolov3_preprocess_func, yolov3_preprocess_func_2, yolov3_variant_preprocess_func, yolov3_variant_preprocess_func_2, yolov3_variant_preprocess_func_3
import onnxruntime
from argparse import Namespace
import os
@ -93,8 +93,10 @@ class YoloV3DataReader(ObejctDetectionDataReader):
def load_serial(self):
width = self.width
height = self.width
nchw_data_list, filename_list, image_size_list = yolov3_preprocess_func(self.image_folder, height, width,
nchw_data_list, filename_list, image_size_list = yolov3_preprocess_func_2(self.image_folder, height, width,
self.start_index, self.stride)
# nchw_data_list, filename_list, image_size_list = yolov3_preprocess_func(self.image_folder, height, width,
# self.start_index, self.stride)
input_name = self.input_name
print("Start from index %s ..." % (str(self.start_index)))
@ -179,18 +181,19 @@ class YoloV3VariantDataReader(YoloV3DataReader):
annotations='./annotations/instances_val2017.json'):
YoloV3DataReader.__init__(self, calibration_image_folder, width, height, start_index, end_index, stride,
batch_size, model_path, is_evaluation, annotations)
self.input_name = '000_net'
# self.input_name = 'images'
# self.input_name = '000_net'
self.input_name = 'images'
def load_serial(self):
width = self.width
height = self.height
input_name = self.input_name
nchw_data_list, filename_list, image_size_list = yolov3_variant_preprocess_func(
# nchw_data_list, filename_list, image_size_list = yolov3_variant_preprocess_func_2(
# self.image_folder, height, width, self.start_index, self.stride)
nchw_data_list, filename_list, image_size_list = yolov3_variant_preprocess_func_3(
self.image_folder, height, width, self.start_index, self.stride)
# nchw_data_list, filename_list, image_size_list = yolov3_variant_2_preprocess_func(
# self.image_folder, height, width, self.start_index, self.stride)
print("Start from index %s ..." % (str(self.start_index)))
data = []
if self.is_evaluation:
img_name_to_img_id = self.img_name_to_img_id

5
onnxruntime/python/tools/quantization/E2E_example_model/object_detection/trt/yolov3/e2e_user_yolov3_example.py
View file

@ -1,5 +1,5 @@
import os
from onnxruntime.quantization import create_calibrator, write_calibration_table
from onnxruntime.quantization import create_calibrator, write_calibration_table, CalibrationMethod
from data_reader import YoloV3DataReader, YoloV3VariantDataReader
from evaluate import YoloV3Evaluator, YoloV3VariantEvaluator
@ -64,7 +64,8 @@ def get_prediction_evaluation(model_path, validation_dataset, providers):
def get_calibration_table_yolov3_variant(model_path, augmented_model_path, calibration_dataset):
calibrator = create_calibrator(model_path, None, augmented_model_path=augmented_model_path)
calibrator = create_calibrator(model_path, [], augmented_model_path=augmented_model_path, calibrate_method=CalibrationMethod.Entropy)
calibrator.set_execution_providers(["CUDAExecutionProvider"])
# DataReader can handle dataset with batch or serial processing depends on its implementation
# Following examples show two different ways to generate calibration table

262
onnxruntime/python/tools/quantization/E2E_example_model/object_detection/trt/yolov3/evaluate.py
View file

@ -10,6 +10,8 @@ import onnxruntime
from onnxruntime.quantization.calibrate import CalibrationDataReader
import numpy as np
import torch
import torchvision
class YoloV3Evaluator:
def __init__(self,
@ -51,6 +53,8 @@ class YoloV3Evaluator:
self.generate_class_to_id(ground_truth_object_class_file)
print(self.class_to_id)
self.session = onnxruntime.InferenceSession(model_path, providers=providers)
def generate_class_to_id(self, ground_truth_object_class_file):
with open(ground_truth_object_class_file) as f:
import json
@ -106,7 +110,7 @@ class YoloV3Evaluator:
})
def predict(self):
session = onnxruntime.InferenceSession(self.model_path, providers=self.providers)
session = self.session
outputs = []
@ -184,23 +188,20 @@ class YoloV3Evaluator:
cocoEval.accumulate()
cocoEval.summarize()
class YoloV3VariantEvaluator(YoloV3Evaluator):
def __init__(self, model_path,
data_reader: CalibrationDataReader,
width=608,
height=384,
providers=["CUDAExecutionProvider"],
ground_truth_object_class_file="./coco-object-categories-2017.json",
onnx_object_class_file="./onnx_coco_classes.txt"):
class YoloV3VariantEvaluator(YoloV3Evaluator):
def __init__(self,
model_path,
data_reader: CalibrationDataReader,
width=608,
height=384,
providers=["CUDAExecutionProvider"],
ground_truth_object_class_file="./coco-object-categories-2017.json",
onnx_object_class_file="./onnx_coco_classes.txt"):
YoloV3Evaluator.__init__(self, model_path, data_reader, width, height, providers,
ground_truth_object_class_file, onnx_object_class_file)
YoloV3Evaluator.__init__(self, model_path, data_reader,width, height, providers, ground_truth_object_class_file, onnx_object_class_file)
def predict(self):
from postprocessing import PostprocessYOLOWrapper
session = onnxruntime.InferenceSession(self.model_path, providers=self.providers)
from postprocessing import PostprocessYOLOWrapper
session = self.session
outputs = []
image_id_list = []
@ -224,24 +225,25 @@ class YoloV3VariantEvaluator(YoloV3Evaluator):
image_size_list = [image_size_list]
image_id_list = [image_id_list]
image_size_batch.append(image_size_list)
image_id_batch.append(image_id_list)
outputs.append(session.run(None, inputs))
for i in range(len(outputs)):
output = outputs[i]
for batch_i in range(self.data_reader.get_batch_size()):
if batch_i > len(image_size_batch[i]) - 1 or batch_i > len(image_id_batch[i]) - 1:
if batch_i > len(image_size_batch[i])-1 or batch_i > len(image_id_batch[i])-1:
continue
image_height = image_size_batch[i][batch_i][0]
image_width = image_size_batch[i][batch_i][1]
image_width= image_size_batch[i][batch_i][1]
image_id = image_id_batch[i][batch_i]
boxes, classes, scores = postprocess_yolo.postprocessor.process(output, (image_width, image_height),
0.01)
boxes, classes, scores = postprocess_yolo.postprocessor.process(
output, (image_width, image_height), 0.01)
for j in range(len(boxes)):
box = boxes[j]
@ -253,13 +255,7 @@ class YoloV3VariantEvaluator(YoloV3Evaluator):
y = float(box[1])
w = float(box[2] - box[0] + 1)
h = float(box[3] - box[1] + 1)
self.prediction_result_list.append({
"image_id": int(image_id),
"category_id": int(id),
"bbox": [x, y, w, h],
"score": scores[j]
})
self.prediction_result_list.append({"image_id":int(image_id), "category_id":int(id), "bbox":[x,y,w,h], "score":scores[j]})
class YoloV3Variant2Evaluator(YoloV3Evaluator):
def __init__(self,
@ -328,7 +324,7 @@ class YoloV3Variant2Evaluator(YoloV3Evaluator):
})
def predict(self):
session = onnxruntime.InferenceSession(self.model_path, providers=self.providers)
session = self.session
outputs = []
image_id_list = []
@ -367,3 +363,213 @@ class YoloV3Variant2Evaluator(YoloV3Evaluator):
image_width = image_size_batch[i][batch_i][1]
image_id = image_id_batch[i][batch_i]
self.set_bbox_prediction(bboxes, scores, image_height, image_width, image_id)
def xywh2xyxy(x):
# Convert nx4 boxes from [x, y, w, h] to [x1, y1, x2, y2] where xy1=top-left, xy2=bottom-right
y = np.zeros_like(x)
y[:, 0] = x[:, 0] - x[:, 2] / 2 # top left x
y[:, 1] = x[:, 1] - x[:, 3] / 2 # top left y
y[:, 2] = x[:, 0] + x[:, 2] / 2 # bottom right x
y[:, 3] = x[:, 1] + x[:, 3] / 2 # bottom right y
return y
def scale_coords(img1_shape, coords, img0_shape, ratio_pad=None):
# Rescale coords (xyxy) from img1_shape to img0_shape
if ratio_pad is None: # calculate from img0_shape
# gain = max(img1_shape) / max(img0_shape) # gain = old / new
gain = min(img1_shape[0] / img0_shape[0], img1_shape[1] / img0_shape[1])
pad = (img1_shape[1] - img0_shape[1] * gain) / 2, (img1_shape[0] - img0_shape[0] * gain) / 2 # wh padding
else:
gain = ratio_pad[0][0]
pad = ratio_pad[1]
coords[:, [0, 2]] -= pad[0] # x padding
coords[:, [1, 3]] -= pad[1] # y padding
coords[:, :4] /= gain
return coords
def letterbox(img, new_shape=(416, 416), color=(114, 114, 114), auto=True, scaleFill=False, scaleup=True):
# Resize image to a 32-pixel-multiple rectangle https://github.com/ultralytics/yolov3/issues/232
shape = img.shape[:2] # current shape [height, width]
if isinstance(new_shape, int):
new_shape = (new_shape, new_shape)
# Scale ratio (new / old)
r = min(new_shape[0] / shape[0], new_shape[1] / shape[1])
if not scaleup: # only scale down, do not scale up (for better test mAP)
r = min(r, 1.0)
# Compute padding
ratio = r, r # width, height ratios
new_unpad = int(round(shape[1] * r)), int(round(shape[0] * r))
dw, dh = new_shape[1] - new_unpad[0], new_shape[0] - new_unpad[1] # wh padding
if auto: # minimum rectangle
dw, dh = np.mod(dw, 64), np.mod(dh, 64) # wh padding
elif scaleFill: # stretch
dw, dh = 0.0, 0.0
new_unpad = new_shape
ratio = new_shape[0] / shape[1], new_shape[1] / shape[0] # width, height ratios
dw /= 2 # divide padding into 2 sides
dh /= 2
if shape[::-1] != new_unpad: # resize
img = cv2.resize(img, new_unpad, interpolation=cv2.INTER_LINEAR)
top, bottom = int(round(dh - 0.1)), int(round(dh + 0.1))
left, right = int(round(dw - 0.1)), int(round(dw + 0.1))
img = cv2.copyMakeBorder(img, top, bottom, left, right, cv2.BORDER_CONSTANT, value=color) # add border
return img, ratio, (dw, dh)
def post_process_without_nms(opts):
final_output = []
for batch_i in range(opt.batch_size):
batch_idx = opts[0][:, 0] == batch_i
bbox = opts[1][batch_idx, :]
score = opts[2][batch_idx, :]
bbox[:, 0] *= opt.input_w #x
bbox[:, 1] *= opt.input_h #y
bbox[:, 2] *= opt.input_w #w
bbox[:, 3] *= opt.input_h #h
bbox = xywh2xyxy(bbox)
bbox0 = scale_coords(img.shape[2:], bbox, img0.shape[0:2])
if bbox0.shape[0] == 0:
final_output.append(torch.empty(0, 5).numpy())
continue
output = np.concatenate((bbox, score), axis=1)
final_output.append(output)
return final_output
def post_process_with_nms(predictions, image_height, image_width, conf_thres=0.35, nms_thres=0.35):
"""Performs NMS and score thresholding
"""
final_output = []
batch_size = 1
input_w = 512
input_h = 288
for batch_i in range(batch_size):
scores = predictions[0][batch_i, :, 0]
keep_idx = scores >= conf_thres
boxes_ = predictions[1][batch_i, keep_idx, :]
boxes_[:, 0] *= input_w #x
boxes_[:, 1] *= input_h #y
boxes_[:, 2] *= input_w #w
boxes_[:, 3] *= input_h #h
boxes_ = xywh2xyxy(boxes_)
img0_shape = (image_height, image_width)
img1_shape = (input_h, input_w)
# bbox = self.scale_coords(img1_shape, bbox, img0_shape)
boxes_ = scale_coords(img1_shape, boxes_, img0_shape)
# boxes_ = scale_coords(img.shape[2:], boxes_, img0.shape[0:2])
boxes_ = torch.from_numpy(boxes_)
scores = torch.from_numpy(scores[keep_idx])
if scores.dim() == 0:
final_output.append(torch.empty(0, 5).numpy())
continue
keep_idx = torchvision.ops.nms(boxes_, scores, nms_thres)
scores = scores[keep_idx].view(-1, 1)
boxes_ = boxes_[keep_idx].view(-1, 4)
output = torch.cat((boxes_, scores), dim=-1)
final_output.append(output.numpy())
return final_output
class YoloV3Variant3Evaluator(YoloV3Evaluator):
def __init__(self,
model_path,
data_reader: CalibrationDataReader,
width=512,
height=288,
providers=["CUDAExecutionProvider"],
ground_truth_object_class_file="./coco-object-categories-2017.json",
onnx_object_class_file="./onnx_coco_classes.txt"):
YoloV3Evaluator.__init__(self, model_path, data_reader, width, height, providers,
ground_truth_object_class_file, onnx_object_class_file)
def set_bbox_prediction(self, bboxes, scores, image_height, image_width, image_id):
for i in range(bboxes.shape[0]):
bbox = bboxes[i]
bbox[0] *= self.width #x
bbox[1] *= self.height #y
bbox[2] *= self.width #w
bbox[3] *= self.height #h
img0_shape = (image_height, image_width)
img1_shape = (self.height, self.width)
bbox = self.xywh2xyxy(bbox)
bbox = self.scale_coords(img1_shape, bbox, img0_shape)
class_name = 'person'
if class_name in self.identical_class_map:
class_name = self.identical_class_map[class_name]
id = self.class_to_id[class_name]
bbox[2] = bbox[2] - bbox[0]
bbox[3] = bbox[3] - bbox[1]
self.prediction_result_list.append({
"image_id": int(image_id),
"category_id": int(id),
"bbox": list(bbox),
"score": scores[i][0]
})
def predict(self):
session = onnxruntime.InferenceSession(self.model_path, providers=self.providers)
outputs = []
image_id_list = []
image_id_batch = []
image_size_list = []
image_size_batch = []
class_name = 'person'
id = self.class_to_id[class_name]
while True:
inputs = self.data_reader.get_next()
if not inputs:
break
image_size_list = inputs["image_size"]
image_id_list = inputs["image_id"]
del inputs["image_size"]
del inputs["image_id"]
# in the case of batch size is 1
if type(image_id_list) == int:
image_size_list = [image_size_list]
image_id_list = [image_id_list]
image_size_batch.append(image_size_list)
image_id_batch.append(image_id_list)
outputs.append(session.run(None, inputs))
for j in range(len(outputs)):
output = outputs[j]
image_id = image_id_batch[j][0]
image_height = image_size_batch[j][0][0]
image_width = image_size_batch[j][0][1]
dets = post_process_with_nms(output, image_height, image_width)[0]
for i in range(dets.shape[0]):
x1 = dets[i, 0]
y1 = dets[i, 1]
x2 = dets[i, 2]
y2 = dets[i, 3]
score = dets[i, 4]
bbox = [x1, y1, x2-x1, y2-y1]
self.prediction_result_list.append({
"image_id": int(image_id),
"category_id": int(id),
"bbox": list(bbox),
"score": score
})

52
onnxruntime/python/tools/quantization/E2E_example_model/object_detection/trt/yolov3/postprocessing.py
View file

@ -1,9 +1,13 @@
import numpy as np
class PostprocessYOLO(object):
"""Class for post-processing the three output tensors from YOLO."""
def __init__(self, yolo_masks, yolo_anchors, nms_threshold, yolo_input_resolution, category_num=80):
def __init__(self,
yolo_masks,
yolo_anchors,
nms_threshold,
yolo_input_resolution,
category_num=80):
"""Initialize with all values that will be kept when processing
several frames. Assuming 3 outputs of the network in the case
of (large) YOLO, or 2 for the Tiny YOLO.
@ -37,7 +41,8 @@ class PostprocessYOLO(object):
for output in outputs:
outputs_reshaped.append(self._reshape_output(output))
boxes_xywh, categories, confidences = self._process_yolo_output(outputs_reshaped, resolution_raw, conf_th)
boxes_xywh, categories, confidences = self._process_yolo_output(
outputs_reshaped, resolution_raw, conf_th)
if len(boxes_xywh) > 0:
# convert (x, y, width, height) to (x1, y1, x2, y2)
@ -46,9 +51,9 @@ class PostprocessYOLO(object):
yy = boxes_xywh[:, 1].reshape(-1, 1)
ww = boxes_xywh[:, 2].reshape(-1, 1)
hh = boxes_xywh[:, 3].reshape(-1, 1)
boxes = np.concatenate([xx, yy, xx + ww, yy + hh], axis=1) + 0.5
boxes[:, [0, 2]] = np.clip(boxes[:, [0, 2]], 0., float(img_w - 1))
boxes[:, [1, 3]] = np.clip(boxes[:, [1, 3]], 0., float(img_h - 1))
boxes = np.concatenate([xx, yy, xx+ww, yy+hh], axis=1) + 0.5
boxes[:, [0, 2]] = np.clip(boxes[:, [0, 2]], 0., float(img_w-1))
boxes[:, [1, 3]] = np.clip(boxes[:, [1, 3]], 0., float(img_h-1))
boxes = boxes.astype(np.int)
else:
boxes = np.zeros((0, 4), dtype=np.int) # empty
@ -118,8 +123,9 @@ class PostprocessYOLO(object):
nscores.append(confidence[keep])
if not nms_categories and not nscores:
return (np.empty((0, 4), dtype=np.float32), np.empty((0, 1),
dtype=np.float32), np.empty((0, 1), dtype=np.float32))
return (np.empty((0, 4), dtype=np.float32),
np.empty((0, 1), dtype=np.float32),
np.empty((0, 1), dtype=np.float32))
boxes = np.concatenate(nms_boxes)
categories = np.concatenate(nms_categories)
@ -136,6 +142,7 @@ class PostprocessYOLO(object):
output_reshaped -- reshaped YOLO output as NumPy arrays with shape (height,width,3,85)
mask -- 2-dimensional tuple with mask specification for this output
"""
def sigmoid_v(array):
return np.reciprocal(np.exp(-array) + 1.0)
@ -238,24 +245,27 @@ class PostprocessYOLO(object):
keep = np.array(keep)
return keep
class PostprocessYOLOWrapper(object):
"""This class encapsulates things needed to run yolo."""
"""Reference from here https://github.com/jkjung-avt/tensorrt_demos/blob/3fb15c908b155d5edc1bf098c6b8c31886cd8e8d/utils/yolo.py"""
def _init_yolov3_postprocessor(self):
h, w = self.input_shape
filters = (self.category_num + 5) * 3
if 'tiny' in self.model:
self.output_shapes = [(1, filters, h // 32, w // 32), (1, filters, h // 16, w // 16)]
self.output_shapes = [(1, filters, h // 32, w // 32),
(1, filters, h // 16, w // 16)]
else:
self.output_shapes = [(1, filters, h // 32, w // 32), (1, filters, h // 16, w // 16),
(1, filters, h // 8, w // 8)]
self.output_shapes = [(1, filters, h // 32, w // 32),
(1, filters, h // 16, w // 16),
(1, filters, h // 8, w // 8)]
if 'tiny' in self.model:
postprocessor_args = {
# A list of 2 three-dimensional tuples for the Tiny YOLO masks
'yolo_masks': [(3, 4, 5), (0, 1, 2)],
# A list of 6 two-dimensional tuples for the Tiny YOLO anchors
'yolo_anchors': [(10, 14), (23, 27), (37, 58), (81, 82), (135, 169), (344, 319)],
'yolo_anchors': [(10, 14), (23, 27), (37, 58),
(81, 82), (135, 169), (344, 319)],
# Threshold for non-max suppression algorithm, float
# value between 0 and 1
'nms_threshold': 0.5,
@ -267,16 +277,14 @@ class PostprocessYOLOWrapper(object):
# A list of 3 three-dimensional tuples for the YOLO masks
'yolo_masks': [(6, 7, 8), (3, 4, 5), (0, 1, 2)],
# A list of 9 two-dimensional tuples for the YOLO anchors
'yolo_anchors': [(10, 13), (16, 30), (33, 23), (30, 61), (62, 45), (59, 119), (116, 90), (156, 198),
(373, 326)],
'yolo_anchors': [(10, 13), (16, 30), (33, 23),
(30, 61), (62, 45), (59, 119),
(116, 90), (156, 198), (373, 326)],
# Threshold for non-max suppression algorithm, float
# value between 0 and 1
'nms_threshold':
0.5,
'yolo_input_resolution':
self.input_shape,
'category_num':
self.category_num
'nms_threshold': 0.5,
'yolo_input_resolution': self.input_shape,
'category_num': self.category_num
}
self.postprocessor = PostprocessYOLO(**postprocessor_args)

143
onnxruntime/python/tools/quantization/E2E_example_model/object_detection/trt/yolov3/preprocessing.py
View file

@ -6,7 +6,6 @@ from PIL import Image
import cv2
import pdb
def yolov3_preprocess_func(images_folder, height, width, start_index=0, size_limit=0):
'''
Loads a batch of images and preprocess them
@ -16,20 +15,19 @@ def yolov3_preprocess_func(images_folder, height, width, start_index=0, size_lim
parameter size_limit: number of images to load. Default is 0 which means all images are picked.
return: list of matrices characterizing multiple images
'''
# this function is from yolo3.utils.letterbox_image
# https://github.com/qqwweee/keras-yolo3/blob/master/yolo3/utils.py
def letterbox_image(image, size):
'''resize image with unchanged aspect ratio using padding'''
iw, ih = image.size
w, h = size
scale = min(w / iw, h / ih)
nw = int(iw * scale)
nh = int(ih * scale)
scale = min(w/iw, h/ih)
nw = int(iw*scale)
nh = int(ih*scale)
image = image.resize((nw, nh), Image.BICUBIC)
new_image = Image.new('RGB', size, (128, 128, 128))
new_image.paste(image, ((w - nw) // 2, (h - nh) // 2))
image = image.resize((nw,nh), Image.BICUBIC)
new_image = Image.new('RGB', size, (128,128,128))
new_image.paste(image, ((w-nw)//2, (h-nh)//2))
return new_image
image_names = os.listdir(images_folder)
@ -44,6 +42,66 @@ def yolov3_preprocess_func(images_folder, height, width, start_index=0, size_lim
else:
batch_filenames = image_names
unconcatenated_batch_data = []
image_size_list = []
print(batch_filenames)
print("size: %s" % str(len(batch_filenames)))
for image_name in batch_filenames:
image_filepath = images_folder + '/' + image_name
img = Image.open(image_filepath)
model_image_size = (height, width)
boxed_image = letterbox_image(img, tuple(reversed(model_image_size)))
image_data = np.array(boxed_image, dtype='float32')
image_data /= 255.
image_data = np.transpose(image_data, [2, 0, 1])
image_data = np.expand_dims(image_data, 0)
unconcatenated_batch_data.append(image_data)
image_size_list.append(np.array([img.size[1], img.size[0]], dtype=np.float32).reshape(1, 2))
batch_data = np.concatenate(np.expand_dims(unconcatenated_batch_data, axis=0), axis=0)
return batch_data, batch_filenames, image_size_list
def yolov3_preprocess_func_2(images_folder, height, width, start_index=0, size_limit=0):
'''
Loads a batch of images and preprocess them
parameter images_folder: path to folder storing images
parameter height: image height in pixels
parameter width: image width in pixels
parameter size_limit: number of images to load. Default is 0 which means all images are picked.
return: list of matrices characterizing multiple images
'''
# reference from here:
# https://github.com/jkjung-avt/tensorrt_demos/blob/3fb15c908b155d5edc1bf098c6b8c31886cd8e8d/utils/yolo.py#L60
def _preprocess_yolo(img, input_shape):
"""Preprocess an image before TRT YOLO inferencing.
# Args
img: int8 numpy array of shape (img_h, img_w, 3)
input_shape: a tuple of (H, W)
# Returns
preprocessed img: float32 numpy array of shape (3, H, W)
"""
img = cv2.resize(img, (input_shape[1], input_shape[0]))
img = cv2.cvtColor(img, cv2.COLOR_BGR2RGB)
img = img.transpose((2, 0, 1)).astype(np.float32)
img /= 255.0
return img
image_names = os.listdir(images_folder)
if start_index >= len(image_names):
return np.asanyarray([]), np.asanyarray([]), np.asanyarray([])
elif size_limit > 0 and len(image_names) >= size_limit:
end_index = start_index + size_limit
if end_index > len(image_names):
end_index = len(image_names)
batch_filenames = [image_names[i] for i in range(start_index, end_index)]
else:
batch_filenames = image_names
unconcatenated_batch_data = []
image_size_list = []
@ -52,7 +110,66 @@ def yolov3_preprocess_func(images_folder, height, width, start_index=0, size_lim
for image_name in batch_filenames:
image_filepath = images_folder + '/' + image_name
img = Image.open(image_filepath)
model_image_size = (height, width)
img = cv2.imread(image_filepath)
image_data = _preprocess_yolo(img, tuple(model_image_size))
image_data = np.ascontiguousarray(image_data)
image_data = np.expand_dims(image_data, 0)
unconcatenated_batch_data.append(image_data)
_height, _width, _ = img.shape
# image_size_list.append(img.shape[0:2]) # img.shape is h, w, c
image_size_list.append(np.array([img.shape[0], img.shape[1]], dtype=np.float32).reshape(1, 2))
batch_data = np.concatenate(np.expand_dims(unconcatenated_batch_data, axis=0), axis=0)
return batch_data, batch_filenames, image_size_list
def yolov3_variant_preprocess_func(images_folder, height, width, start_index=0, size_limit=0):
'''
Loads a batch of images and preprocess them
parameter images_folder: path to folder storing images
parameter height: image height in pixels
parameter width: image width in pixels
parameter size_limit: number of images to load. Default is 0 which means all images are picked.
return: list of matrices characterizing multiple images
'''
# this function is from yolo3.utils.letterbox_image
# https://github.com/qqwweee/keras-yolo3/blob/master/yolo3/utils.py
def letterbox_image(image, size):
'''resize image with unchanged aspect ratio using padding'''
iw, ih = image.size
w, h = size
scale = min(w/iw, h/ih)
nw = int(iw*scale)
nh = int(ih*scale)
image = image.resize((nw,nh), Image.BICUBIC)
new_image = Image.new('RGB', size, (128,128,128))
new_image.paste(image, ((w-nw)//2, (h-nh)//2))
return new_image
image_names = os.listdir(images_folder)
if start_index >= len(image_names):
return np.asanyarray([]), np.asanyarray([]), np.asanyarray([])
elif size_limit > 0 and len(image_names) >= size_limit:
end_index = start_index + size_limit
if end_index > len(image_names):
end_index = len(image_names)
batch_filenames = [image_names[i] for i in range(start_index, end_index)]
else:
batch_filenames = image_names
unconcatenated_batch_data = []
image_size_list = []
print(batch_filenames)
print("size: %s" % str(len(batch_filenames)))
for image_name in batch_filenames:
image_filepath = images_folder + '/' + image_name
img = Image.open(image_filepath)
model_image_size = (height, width)
boxed_image = letterbox_image(img, tuple(reversed(model_image_size)))
image_data = np.array(boxed_image, dtype='float32')
@ -60,13 +177,13 @@ def yolov3_preprocess_func(images_folder, height, width, start_index=0, size_lim
image_data = np.transpose(image_data, [2, 0, 1])
image_data = np.expand_dims(image_data, 0)
unconcatenated_batch_data.append(image_data)
image_size_list.append(np.array([img.size[1], img.size[0]], dtype=np.float32).reshape(1, 2))
image_size_list.append((img.size[1], img.size[0])) # img.shape is h, w, c
# image_size_list.append(np.array([img.size[1], img.size[0]], dtype=np.float32).reshape(1, 2))
batch_data = np.concatenate(np.expand_dims(unconcatenated_batch_data, axis=0), axis=0)
return batch_data, batch_filenames, image_size_list
def yolov3_variant_preprocess_func(images_folder, height, width, start_index=0, size_limit=0):
def yolov3_variant_preprocess_func_2(images_folder, height, width, start_index=0, size_limit=0):
'''
Loads a batch of images and preprocess them
parameter images_folder: path to folder storing images
@ -127,7 +244,7 @@ def yolov3_variant_preprocess_func(images_folder, height, width, start_index=0,
# This is for special tuned yolov3 model
def yolov3_variant_2_preprocess_func(images_folder, height, width, start_index=0, size_limit=0):
def yolov3_variant_preprocess_func_3(images_folder, height, width, start_index=0, size_limit=0):
def letterbox(img, new_shape=(416, 416), color=(114, 114, 114), auto=True, scaleFill=False, scaleup=True):
# Resize image to a 32-pixel-multiple rectangle https://github.com/ultralytics/yolov3/issues/232
shape = img.shape[:2] # current shape [height, width]

2
onnxruntime/python/tools/quantization/__init__.py
View file

@ -1,4 +1,4 @@
from .quantize import quantize, quantize_static, quantize_dynamic, quantize_qat
from .quantize import QuantizationMode
from .calibrate import CalibrationDataReader, CalibraterBase, MinMaxCalibrater, create_calibrator
from .calibrate import CalibrationDataReader, CalibraterBase, MinMaxCalibrater, create_calibrator, CalibrationMethod
from .quant_utils import QuantType, QuantFormat, write_calibration_table

304
onnxruntime/python/tools/quantization/calibrate.py
View file

@ -10,12 +10,12 @@ import os
import numpy as np
import onnx
import onnxruntime
from onnx import helper, TensorProto, ModelProto
from onnx import helper, TensorProto, ModelProto, shape_inference
from onnx import onnx_pb as onnx_proto
from six import string_types
from enum import Enum
from .quant_utils import QuantType
from .quant_utils import QuantType, smooth_distribution
from .registry import QLinearOpsRegistry
import abc
@ -24,6 +24,7 @@ import itertools
class CalibrationMethod(Enum):
MinMax = 0
Entropy = 1
class CalibrationDataReader(metaclass=abc.ABCMeta):
@ -51,6 +52,9 @@ class CalibraterBase:
else:
raise ValueError('model should be either model path or onnx.ModelProto.')
# Apply shape inference on the model
self.model = onnx.shape_inference.infer_shapes(self.model)
self.op_types_to_calibrate = op_types_to_calibrate
self.augmented_model_path = augmented_model_path
@ -80,6 +84,33 @@ class CalibraterBase:
sess_options=sess_options,
providers=self.execution_providers)
def select_tensors_to_calibrate(self, model):
'''
select all quantization_candidates op type nodes' input/output tensors.
returns:
tensors (set): set of tensor name.
value_infos (dict): tensor name to value info.
'''
value_infos = {vi.name: vi for vi in model.graph.value_info}
value_infos.update({ot.name: ot for ot in model.graph.output})
value_infos.update({it.name: it for it in model.graph.input})
initializer = set(init.name for init in model.graph.initializer)
tensors_to_calibrate = set()
tensor_type_to_calibrate = set([TensorProto.FLOAT, TensorProto.FLOAT16])
for node in model.graph.node:
if len(self.op_types_to_calibrate) == 0 or node.op_type in self.op_types_to_calibrate:
for tensor_name in itertools.chain(node.input, node.output):
if tensor_name in value_infos.keys():
vi = value_infos[tensor_name]
if vi.type.HasField('tensor_type') and (
vi.type.tensor_type.elem_type in tensor_type_to_calibrate) and (
tensor_name not in initializer):
tensors_to_calibrate.add(tensor_name)
return tensors_to_calibrate, value_infos
def get_augment_model(self):
'''
return: augmented onnx model
@ -129,44 +160,30 @@ class MinMaxCalibrater(CalibraterBase):
model = onnx_proto.ModelProto()
model.CopyFrom(self.model)
model = onnx.shape_inference.infer_shapes(model)
value_infos = {vi.name: vi for vi in model.graph.value_info}
value_infos.update({ot.name: ot for ot in model.graph.output})
value_infos.update({it.name: it for it in model.graph.input})
initializer = set(init.name for init in model.graph.initializer)
added_nodes = []
added_outputs = []
tensors_to_calibrate = set()
tensor_type_to_calibrate = set([TensorProto.FLOAT, TensorProto.FLOAT16])
tensors, value_infos = self.select_tensors_to_calibrate(model)
for node in model.graph.node:
if len(self.op_types_to_calibrate) == 0 or node.op_type in self.op_types_to_calibrate:
for tensor_name in itertools.chain(node.input, node.output):
if tensor_name in value_infos.keys():
vi = value_infos[tensor_name]
if vi.type.HasField('tensor_type') and (
vi.type.tensor_type.elem_type in tensor_type_to_calibrate) and (
tensor_name not in initializer):
tensors_to_calibrate.add(tensor_name)
for tensor in tensors:
# Get tensor's shape
dim = len(value_infos[tensor].type.tensor_type.shape.dim)
shape = (1,) if dim == 1 else list(1 for i in range(dim))
for tensor in tensors_to_calibrate:
# Adding ReduceMin nodes
reduce_min_name = tensor + '_ReduceMin'
reduce_min_node = onnx.helper.make_node('ReduceMin', [tensor], [tensor + '_ReduceMin'],
reduce_min_name,
keepdims=0)
reduce_min_node = onnx.helper.make_node('ReduceMin', [tensor], [tensor + '_ReduceMin'], reduce_min_name)
added_nodes.append(reduce_min_node)
added_outputs.append(helper.make_tensor_value_info(reduce_min_node.output[0], TensorProto.FLOAT, ()))
added_outputs.append(helper.make_tensor_value_info(reduce_min_node.output[0], TensorProto.FLOAT, shape))
# Adding ReduceMax nodes
reduce_max_name = tensor + '_ReduceMax'
reduce_max_node = onnx.helper.make_node('ReduceMax', [tensor], [tensor + '_ReduceMax'],
reduce_max_name,
keepdims=0)
reduce_max_node = onnx.helper.make_node('ReduceMax', [tensor], [tensor + '_ReduceMax'], reduce_max_name)
added_nodes.append(reduce_max_node)
added_outputs.append(helper.make_tensor_value_info(reduce_max_node.output[0], TensorProto.FLOAT, ()))
added_outputs.append(helper.make_tensor_value_info(reduce_max_node.output[0], TensorProto.FLOAT, shape))
model.graph.node.extend(added_nodes)
model.graph.output.extend(added_outputs)
@ -226,6 +243,239 @@ class MinMaxCalibrater(CalibraterBase):
return self.calibrate_tensors_range
class EntropyCalibrater(CalibraterBase):
def __init__(self, model, op_types_to_calibrate=[], augmented_model_path='augmented_model.onnx'):
'''
:param model: ONNX model to calibrate. It can be a ModelProto or a model path
:param op_types_to_calibrate: operator types to calibrate. By default, calibrate all the float32/float16 tensors.
:param augmented_model_path: save augmented model to this path.
'''
super(EntropyCalibrater, self).__init__(model, op_types_to_calibrate, augmented_model_path)
self.intermediate_outputs = []
self.calibrate_tensors_range = None
self.num_model_outputs = len(self.model.graph.output)
self.model_original_outputs = set(output.name for output in self.model.graph.output)
self.collector = None
def augment_graph(self):
'''
make all quantization_candidates op type nodes as part of the graph output.
:return: augmented ONNX model
'''
model = onnx_proto.ModelProto()
model.CopyFrom(self.model)
model = onnx.shape_inference.infer_shapes(model)
added_nodes = []
added_outputs = []
tensors, value_infos = self.select_tensors_to_calibrate(model)
for tensor in tensors:
added_outputs.append(value_infos[tensor])
model.graph.node.extend(added_nodes)
model.graph.output.extend(added_outputs)
onnx.save(model, self.augmented_model_path)
self.augment_model = model
def clear_collected_data(self):
self.intermediate_outputs = []
def collect_data(self, data_reader: CalibrationDataReader):
'''
Entropy Calibrator collects operators' tensors as well as generates tensor histogram for each operator.
'''
while True:
inputs = data_reader.get_next()
if not inputs:
break
self.intermediate_outputs.append(self.infer_session.run(None, inputs))
if len(self.intermediate_outputs) == 0:
raise ValueError("No data is collected.")
output_names = [self.infer_session.get_outputs()[i].name for i in range(len(self.intermediate_outputs[0]))]
output_dicts_list = [
dict(zip(output_names, intermediate_output)) for intermediate_output in self.intermediate_outputs
]
merged_dict = {}
for d in output_dicts_list:
for k, v in d.items():
merged_dict.setdefault(k, []).append(v)
clean_merged_dict = dict((i, merged_dict[i]) for i in merged_dict if i not in self.model_original_outputs)
if not self.collector:
self.collector = HistogramCollector()
self.collector.collect(clean_merged_dict)
def compute_range(self):
'''
Compute the min-max range of tensor
:return: dictionary mapping: {added node names: (ReduceMin, ReduceMax) pairs }
'''
if not self.collector:
raise ValueError("No collector created and can't generate calibration data.")
return self.collector.get_optimal_collection_result()
class CalibrationDataCollector(metaclass=abc.ABCMeta):
"""
Base class for collecting data for calibration-based quantization.
"""
@abc.abstractmethod
def collect(self, name_to_arr):
"""
Generate informative data based on given data.
name_to_arr : dict
tensor name to NDArray data
"""
raise NotImplementedError
@abc.abstractmethod
def get_optimal_collection_result(self):
"""
Get the optimal result among collection data.
"""
raise NotImplementedError
class HistogramCollector(CalibrationDataCollector):
"""
Implementation of collecting histogram data as dict for each tensor targeting on entropy calibration.
ref: https://github.com//apache/incubator-mxnet/blob/master/python/mxnet/contrib/quantization.py
"""
def __init__(self, num_quantized_bins=128):
self.histogram_dict = {}
self.num_quantized_bins= num_quantized_bins
def get_histogram_dict(self):
return self.histogram_dict
def collect(self, name_to_arr):
for tensor, data_arr in name_to_arr.items():
data_arr = np.asarray(data_arr)
data_arr = data_arr.flatten()
if data_arr.size > 0:
min_value = np.min(data_arr)
max_value = np.max(data_arr)
else:
min_value = 0
max_value = 0
threshold = max(abs(min_value), abs(max_value))
if tensor in self.histogram_dict:
old_histogram = self.histogram_dict[tensor]
self.histogram_dict[tensor] = self.merge_histogram(old_histogram, data_arr, min_value, max_value, threshold)
else:
# hist, hist_edges = np.histogram(data_arr, self.num_quantized_bins, range=(min_value, max_value))
hist, hist_edges = np.histogram(data_arr, self.num_quantized_bins, range=(-threshold, threshold))
self.histogram_dict[tensor] = (hist, hist_edges, min_value, max_value, threshold)
def merge_histogram(self, old_histogram, data_arr, new_min, new_max, new_threshold):
(old_hist, old_hist_edges, old_min, old_max, old_threshold) = old_histogram
if new_threshold <= old_threshold:
new_hist, _ = np.histogram(data_arr, len(old_hist), range=(-old_threshold, old_threshold))
return (new_hist + old_hist, old_hist_edges, min(old_min, new_min), max(old_max, new_max), old_threshold)
else:
if old_threshold == 0:
hist, hist_edges = np.histogram(data_arr, new_num_bins, range=(-new_threshold, new_threshold))
hist[len(hist) // 2] += len(old_hist)
else:
old_num_bins = len(old_hist)
old_stride = 2 * old_threshold / old_num_bins
half_increased_bins = int((new_threshold - old_threshold) // old_stride + 1)
new_num_bins = old_num_bins + 2 * half_increased_bins
new_threshold = half_increased_bins * old_stride + old_threshold
hist, hist_edges = np.histogram(data_arr, new_num_bins, range=(-new_threshold, new_threshold))
hist[half_increased_bins:new_num_bins-half_increased_bins] += old_hist
return (hist, hist_edges, min(old_min, new_min), max(old_max, new_max), new_threshold)
def get_optimal_collection_result(self):
histogram_dict = self.histogram_dict
num_quantized_bins = self.num_quantized_bins
thresholds_dict = {} # per tensor thresholds
for tensor, histogram in histogram_dict.items():
optimal_threshold = self.get_optimal_threshold(histogram, num_quantized_bins)
thresholds_dict[tensor] = optimal_threshold
return thresholds_dict
def get_optimal_threshold(self, histogram, num_quantized_bins):
from scipy.stats import entropy
import copy
hist, hist_edges, _, _, _ = histogram
num_bins = hist.size
zero_bin_index = num_bins // 2
num_half_quantized_bin = num_quantized_bins // 2
kl_divergence = np.zeros(zero_bin_index - num_half_quantized_bin + 1)
thresholds = [(0, 0) for i in range(kl_divergence.size)]
for i in range(num_half_quantized_bin, zero_bin_index + 1, 1):
start_index = zero_bin_index - i
end_index = zero_bin_index + i + 1 if (zero_bin_index + i + 1) <= num_bins else num_bins
thresholds[i - num_half_quantized_bin] = (float(hist_edges[start_index]), float(hist_edges[end_index]))
sliced_distribution = copy.deepcopy(hist[start_index:end_index])
# reference distribution p
p = sliced_distribution.copy() # a copy of np array
left_outliers_count = sum(hist[:start_index])
right_outliers_count = sum(hist[end_index:])
p[0] += left_outliers_count
p[-1] += right_outliers_count
# nonzeros[i] incidates whether p[i] is non-zero
nonzeros = (p != 0).astype(np.int64)
# quantize p.size bins into quantized bins (default 128 bins)
quantized_bins = np.zeros(num_quantized_bins, dtype=np.int64)
num_merged_bins = sliced_distribution.size // num_quantized_bins
# merge bins into quantized bins
for index in range(num_quantized_bins):
start = index * num_merged_bins
end = start + num_merged_bins
quantized_bins[index] = sum(sliced_distribution[start:end])
quantized_bins[-1] += sum(sliced_distribution[num_quantized_bins * num_merged_bins:])
# in order to compare p and q, we need to make length of q equals to length of p
# expand quantized bins into p.size bins
q = np.zeros(p.size, dtype=np.int64)
for index in range(num_quantized_bins):
start = index * num_merged_bins
end = start + num_merged_bins
norm = sum(nonzeros[start:end])
if norm != 0:
q[start:end] = float(quantized_bins[index]) / float(norm)
p = smooth_distribution(p)
q = smooth_distribution(q)
if isinstance(q, np.ndarray):
kl_divergence[i - num_half_quantized_bin] = entropy(p, q)
else:
kl_divergence[i - num_half_quantized_bin] = float('inf')
min_kl_divergence_idx = np.argmin(kl_divergence)
optimal_threshold = thresholds[min_kl_divergence_idx]
return optimal_threshold
def create_calibrator(model,
op_types_to_calibrate=[],
@ -233,5 +483,7 @@ def create_calibrator(model,
calibrate_method=CalibrationMethod.MinMax):
if calibrate_method == CalibrationMethod.MinMax:
return MinMaxCalibrater(model, op_types_to_calibrate, augmented_model_path)
elif calibrate_method == CalibrationMethod.Entropy:
return EntropyCalibrater(model, op_types_to_calibrate, augmented_model_path)
raise ValueError('Unsupported calibration method {}'.format(calibrate_method))

5
onnxruntime/python/tools/quantization/onnx_model.py
View file

@ -232,14 +232,15 @@ class ONNXModel:
unused_nodes = []
nodes = self.nodes()
for node in nodes:
if node.op_type == "Constant" and node.output[0] not in input_name_to_nodes:
if node.op_type == "Constant" and not self.is_graph_output(
node.output[0]) and node.output[0] not in input_name_to_nodes:
unused_nodes.append(node)
self.remove_nodes(unused_nodes)
ununsed_weights = []
for w in self.initializer():
if w.name not in input_name_to_nodes:
if w.name not in input_name_to_nodes and not self.is_graph_output(w.name):
ununsed_weights.append(w)
# Remove from graph.input
for graph_input in self.graph().input:

11
onnxruntime/python/tools/quantization/onnx_quantizer.py
View file

@ -531,6 +531,10 @@ class ONNXQuantizer:
Quantized the bias. Zero Point == 0 and Scale == Input_Scale * Weight_Scale
'''
# Handle case where bias already in quantizatio map
if bias_name in self.quantized_value_map:
return self.quantized_value_map[bias_name].q_name
# get scale for weight
weight_scale_name = self.quantized_value_map[weight_name].scale_name
weight_initializer = find_by_name(weight_scale_name, self.model.initializer())
@ -571,8 +575,7 @@ class ONNXQuantizer:
assert (bias_name not in self.quantized_value_map)
quantized_value = QuantizedValue(bias_name, quantized_bias_name, quantized_bias_scale_name, "",
QuantizedValueType.Initializer, 0 if bias_scale_data.size > 1 else None,
onnx_proto.TensorProto.INT32)
QuantizedValueType.Initializer, 0 if bias_scale_data.size > 1 else None)
self.quantized_value_map[bias_name] = quantized_value
return quantized_bias_name
@ -664,7 +667,7 @@ class ONNXQuantizer:
# Log entry for this quantized weight
quantized_value = QuantizedValue(weight.name, q_weight_name, scale_name, zp_name,
QuantizedValueType.Initializer, None, qType)
QuantizedValueType.Initializer, None)
self.quantized_value_map[weight.name] = quantized_value
return q_weight_name, zp_name, scale_name
@ -710,7 +713,7 @@ class ONNXQuantizer:
scale_name = weight_name + "_scale"
quantized_value = QuantizedValue(weight_name, q_weight_name, scale_name, zp_name,
QuantizedValueType.Initializer, None, weight_qType)
QuantizedValueType.Initializer, None)
self.quantized_value_map[weight_name] = quantized_value
# Update packed weight, zero point, and scale initializers

2
onnxruntime/python/tools/quantization/operators/gavgpool.py
View file

@ -24,7 +24,7 @@ class QGlobalAveragePool(QuantOperatorBase):
output_scale_name = output_scale_name_from_parameter if data_found else quantized_input_value.scale_name
output_zp_name = output_zp_name_from_parameter if data_found else quantized_input_value.zp_name
quantized_output_value = QuantizedValue(node.output[0], node.output[0] + "_quantized", output_scale_name,
output_zp_name, quantized_input_value.qType)
output_zp_name)
self.quantizer.quantized_value_map[node.output[0]] = quantized_output_value
kwargs = {}

4
onnxruntime/python/tools/quantization/operators/pad.py
View file

@ -39,7 +39,7 @@ class QPad(QuantOperatorBase):
scale_array = onnx.numpy_helper.to_array(scale_tensor)
scale_value = scale_array.item() if scale_array.ndim == 0 else scale_array[0]
padding_constant_array = onnx.numpy_helper.to_array(padding_constant_initializer)
quantized_padding_constant_array = quantize_nparray(quantized_input_value.qType,
quantized_padding_constant_array = quantize_nparray(self.quantizer.input_qType,
padding_constant_array, scale_value, zp_value)
quantized_padding_constant_name = node.input[2] + "_quantized"
quantized_padding_constant_initializer = onnx.numpy_helper.from_array(
@ -49,7 +49,7 @@ class QPad(QuantOperatorBase):
self.quantizer.model.add_initializer(quantized_padding_constant_initializer)
node.input[2] = quantized_padding_constant_name
else:
pad_value_qnodes = self.quantizer._get_quantize_input_nodes(node, 2, quantized_input_value.qType,
pad_value_qnodes = self.quantizer._get_quantize_input_nodes(node, 2, self.quantizer.input_qType,
quantized_input_value.scale_name,
quantized_input_value.zp_name)
self.quantizer.new_nodes += [pad_value_qnodes]

31
onnxruntime/python/tools/quantization/operators/reshape.py Normal file
View file

@ -0,0 +1,31 @@
import onnx
from .base_operator import QuantOperatorBase
from ..quant_utils import QuantizedValue, QuantizedValueType
from onnx import onnx_pb as onnx_proto
class ReshapeQuant(QuantOperatorBase):
def __init__(self, onnx_quantizer, onnx_node):
super().__init__(onnx_quantizer, onnx_node)
def quantize(self):
node = self.node
assert (node.op_type == "Reshape")
# If input to this node is not quantized then keep this node
if node.input[0] not in self.quantizer.quantized_value_map:
self.quantizer.new_nodes += [node]
return
# Reshape is a no-op in terms of quantization
quantized_input_value = self.quantizer.quantized_value_map[node.input[0]]
quantized_output_value = QuantizedValue(node.output[0], node.output[0] + "_quantized",
quantized_input_value.scale_name, quantized_input_value.zp_name,
QuantizedValueType.Input)
# Create an entry for output quantized value
self.quantizer.quantized_value_map[node.output[0]] = quantized_output_value
node.input[0] = quantized_input_value.q_name
node.output[0] = quantized_output_value.q_name
self.quantizer.new_nodes += [node]

2
onnxruntime/python/tools/quantization/qdq_quantizer.py
View file

@ -134,7 +134,7 @@ class QDQQuantizer(ONNXQuantizer):
self.model.add_nodes([qlinear_node, dequant_node])
quantized_value = QuantizedValue(tensor_name, tensor_name + "_QuantizeLinear", scale_name, zp_name,
QuantizedValueType.Input, None, self.input_qType)
QuantizedValueType.Input)
self.quantized_value_map[tensor_name] = quantized_value
def quantize_bias_tensors(self):

55
onnxruntime/python/tools/quantization/quant_utils.py
View file

@ -93,15 +93,15 @@ class QuantFormat(Enum):
except KeyError:
raise ValueError()
QUANT_TYPE_TO_NP_TYPE = {
QuantType.QInt8: numpy.dtype('int8'),
QuantType.QUInt8: numpy.dtype('uint8'),
ONNX_TYPE_TO_NP_TYPE = {
onnx_proto.TensorProto.INT8: numpy.dtype('int8'),
onnx_proto.TensorProto.UINT8: numpy.dtype('uint8')
}
def quantize_nparray(qtype, arr, scale, zero_point, low=None, high=None):
dtype = QUANT_TYPE_TO_NP_TYPE[qtype]
def quantize_nparray(qType, arr, scale, zero_point, low=None, high=None):
assert qType in ONNX_TYPE_TO_NP_TYPE, \
"Unexpected data type {} requested. Only INT8 and UINT8 are supported.".format(qType)
dtype = ONNX_TYPE_TO_NP_TYPE[qType]
cliplow = max(0 if dtype == numpy.uint8 else -127, -127 if low is None else low)
cliphigh = min(255 if dtype == numpy.uint8 else 127, 255 if high is None else high)
arr_fp32 = numpy.asarray((arr.astype(numpy.float32) / scale).round() + zero_point)
@ -144,12 +144,7 @@ def quantize_data(data, quantize_range, qType):
rmax = max(max(data), 0)
zero_point, scale = compute_scale_zp(rmin, rmax, qType, quantize_range)
if qType == onnx_proto.TensorProto.INT8:
quantized_data = quantize_nparray(QuantType.QInt8, numpy.asarray(data), scale, zero_point)
elif qType == onnx_proto.TensorProto.UINT8:
quantized_data = quantize_nparray(QuantType.QUInt8, numpy.asarray(data), scale, zero_point)
else:
raise ValueError("Unexpected data type {} requested. Only INT8 and UINT8 are supported.".format(qType))
quantized_data = quantize_nparray(qType, numpy.asarray(data), scale, zero_point)
return rmin, rmax, zero_point, scale, quantized_data
@ -181,8 +176,7 @@ class QuantizedInitializer:
scales,
data=[],
quantized_data=[],
axis=None,
qType=QuantType.QUInt8):
axis=None):
self.name = name
self.initializer = initializer # TensorProto initializer in ONNX graph
self.rmins = rmins # List of minimum range for each axis
@ -195,7 +189,6 @@ class QuantizedInitializer:
# Scalar to specify which dimension in the initializer to weight pack.
self.axis = axis
# If empty, single zero point and scales computed from a single rmin and rmax
self.qType = qType # type of quantized data.
class QuantizedValue:
@ -208,15 +201,13 @@ class QuantizedValue:
scale_name,
zero_point_name,
quantized_value_type,
axis=None,
qType=QuantType.QUInt8):
axis=None):
self.original_name = name
self.q_name = new_quantized_name
self.scale_name = scale_name
self.zp_name = zero_point_name
self.value_type = quantized_value_type
self.axis = axis
self.qType = qType
class BiasToQuantize:
@ -368,3 +359,29 @@ def write_calibration_table(calibration_cache):
s = key + ' ' + str(max(abs(value[0]), abs(value[1])))
file.write(s)
file.write('\n')
def smooth_distribution(p, eps=0.0001):
"""Given a discrete distribution (may have not been normalized to 1),
smooth it by replacing zeros with eps multiplied by a scaling factor
and taking the corresponding amount off the non-zero values.
Ref: http://web.engr.illinois.edu/~hanj/cs412/bk3/KL-divergence.pdf
https://github.com//apache/incubator-mxnet/blob/master/python/mxnet/contrib/quantization.py
"""
import numpy as np
is_zeros = (p == 0).astype(np.float32)
is_nonzeros = (p != 0).astype(np.float32)
n_zeros = is_zeros.sum()
n_nonzeros = p.size - n_zeros
if not n_nonzeros:
# raise ValueError('The discrete probability distribution is malformed. All entries are 0.')
return -1
eps1 = eps * float(n_zeros) / float(n_nonzeros)
assert eps1 < 1.0, 'n_zeros=%d, n_nonzeros=%d, eps1=%f' % (n_zeros, n_nonzeros, eps1)
hist = p.astype(np.float32)
hist += eps * is_zeros + (-eps1) * is_nonzeros
assert (hist <= 0).sum() == 0
return hist

11
onnxruntime/python/tools/quantization/quantize.py
View file

@ -24,7 +24,7 @@ from .registry import QLinearOpsRegistry, IntegerOpsRegistry
from .onnx_model import ONNXModel
from .onnx_quantizer import ONNXQuantizer
from .qdq_quantizer import QDQQuantizer
from .calibrate import CalibrationDataReader, create_calibrator
from .calibrate import CalibrationDataReader, create_calibrator, CalibrationMethod
def optimize_model(model_path: Path):
@ -145,7 +145,9 @@ def quantize_static(model_input,
nodes_to_quantize=[],
nodes_to_exclude=[],
optimize_model=True,
use_external_data_format=False):
use_external_data_format=False,
calibrate_method=CalibrationMethod.MinMax):
'''
Given an onnx model and calibration data reader, create a quantized onnx model and save it into a file
:param model_input: file path of model to quantize
@ -173,6 +175,9 @@ def quantize_static(model_input,
when it is not None.
:param optimize_model: optimize model before quantization.
:parma use_external_data_format: option used for large size (>2GB) model. Set to False by default.
:param calibrate_method:
Current calibration methods supported are MinMax and Entropy.
Please use CalibrationMethod.MinMax or CalibrationMethod.Entropy as options.
'''
if activation_type != QuantType.QUInt8:
@ -185,7 +190,7 @@ def quantize_static(model_input,
model = load_model(Path(model_input), optimize_model)
calibrator = create_calibrator(model, op_types_to_quantize)
calibrator = create_calibrator(model, op_types_to_quantize, calibrate_method=calibrate_method)
calibrator.collect_data(calibration_data_reader)
tensors_range = calibrator.compute_range()

5
onnxruntime/python/tools/quantization/registry.py
View file

@ -13,8 +13,11 @@ from .operators.gavgpool import QGlobalAveragePool
from .operators.lstm import LSTMQuant
from .operators.split import QSplit
from .operators.pad import QPad
from .operators.reshape import ReshapeQuant
CommonOpsRegistry = {"Gather": GatherQuant, "EmbedLayerNormalization": EmbedLayerNormalizationQuant}
CommonOpsRegistry = {"Gather": GatherQuant,
"EmbedLayerNormalization": EmbedLayerNormalizationQuant,
"Reshape": ReshapeQuant}
IntegerOpsRegistry = {
"Conv": ConvInteger,

206
onnxruntime/python/tools/tensorrt/perf/benchmark.py
View file

@ -235,72 +235,6 @@ def get_acl_version():
version = version_match.group(0).split(' ')[0]
return version
def get_cuda_version():
from pathlib import Path
home = str(Path.home())
p1 = subprocess.Popen(["find", home+"/.local/lib/", "-name", "onnxruntime_pybind11_state.so"], stdout=subprocess.PIPE)
stdout, sterr = p1.communicate()
stdout = stdout.decode("ascii").strip()
p1 = subprocess.Popen(["ldd", stdout], stdout=subprocess.PIPE)
p2 = subprocess.Popen(["grep", "libcudart.so"], stdin=p1.stdout, stdout=subprocess.PIPE)
stdout, sterr = p2.communicate()
stdout = stdout.decode("ascii").strip()
return stdout
def get_trt_version():
from pathlib import Path
home = str(Path.home())
p1 = subprocess.Popen(["find", home+"/.local/lib/", "-name", "onnxruntime_pybind11_state.so"], stdout=subprocess.PIPE)
stdout, sterr = p1.communicate()
stdout = stdout.decode("ascii").strip()
p1 = subprocess.Popen(["ldd", stdout], stdout=subprocess.PIPE)
p2 = subprocess.Popen(["grep", "libnvinfer.so"], stdin=p1.stdout, stdout=subprocess.PIPE)
stdout, sterr = p2.communicate()
stdout = stdout.decode("ascii").strip()
if stdout == "":
p1 = subprocess.Popen(["find", home+"/.local/lib/", "-name", "libonnxruntime_providers_tensorrt.so"], stdout=subprocess.PIPE)
stdout, sterr = p1.communicate()
stdout = stdout.decode("ascii").strip()
p1 = subprocess.Popen(["ldd", stdout], stdout=subprocess.PIPE)
p2 = subprocess.Popen(["grep", "libnvinfer.so"], stdin=p1.stdout, stdout=subprocess.PIPE)
stdout, sterr = p2.communicate()
stdout = stdout.decode("ascii").strip()
return stdout
# not use for this script temporarily
def tmp_get_trt_version():
p1 = subprocess.Popen(["dpkg", "-l"], stdout=subprocess.PIPE)
p2 = subprocess.Popen(["grep", "TensorRT runtime libraries"], stdin=p1.stdout, stdout=subprocess.PIPE)
stdout, sterr = p2.communicate()
stdout = stdout.decode("ascii").strip()
if stdout != "":
stdout = re.sub('\s+', ' ', stdout)
return stdout
if os.path.exists("/usr/lib/x86_64-linux-gnu/libnvinfer.so"):
p1 = subprocess.Popen(["readelf", "-s", "/usr/lib/x86_64-linux-gnu/libnvinfer.so"], stdout=subprocess.PIPE)
p2 = subprocess.Popen(["grep", "version"], stdin=p1.stdout, stdout=subprocess.PIPE)
stdout, sterr = p2.communicate()
stdout = stdout.decode("ascii").strip()
stdout = stdout.split(" ")[-1]
return stdout
elif os.path.exists("/usr/lib/aarch64-linux-gnu/libnvinfer.so"):
p1 = subprocess.Popen(["readelf", "-s", "/usr/lib/aarch64-linux-gnu/libnvinfer.so"], stdout=subprocess.PIPE)
p2 = subprocess.Popen(["grep", "version"], stdin=p1.stdout, stdout=subprocess.PIPE)
stdout, sterr = p2.communicate()
stdout = stdout.decode("ascii").strip()
stdout = stdout.split(" ")[-1]
return stdout
return ""
#######################################################################################################################################
# The following two lists will be generated.
#
@ -309,7 +243,6 @@ def tmp_get_trt_version():
#######################################################################################################################################
def load_onnx_model_zoo_test_data(path, all_inputs_shape, data_type="fp32"):
logger.info("Parsing test data in {} ...".format(path))
# p1 = subprocess.Popen(["find", path, "-name", "test_data_set*", "-type", "d"], stdout=subprocess.PIPE)
p1 = subprocess.Popen(["find", path, "-name", "test_data*", "-type", "d"], stdout=subprocess.PIPE)
p2 = subprocess.Popen(["sort"], stdin=p1.stdout, stdout=subprocess.PIPE)
stdout, sterr = p2.communicate()
@ -419,7 +352,15 @@ def generate_onnx_model_random_input(test_times, ref_input):
return inputs
def validate(all_ref_outputs, all_outputs, rtol=0, atol=1.5):
def percentage_in_allowed_threshold(e, percent_mismatch):
percent_string = re.search(r'\(([^)]+)', str(e)).group(1)
if "%" in percent_string:
percentage_wrong = float(percent_string.replace("%",""))
return percentage_wrong < percent_mismatch
else:
return False # error in output
def validate(all_ref_outputs, all_outputs, rtol, atol, percent_mismatch):
if len(all_ref_outputs) == 0:
logger.info("No reference output provided.")
return True, None
@ -428,22 +369,25 @@ def validate(all_ref_outputs, all_outputs, rtol=0, atol=1.5):
logger.info('Predicted {} results.'.format(len(all_outputs)))
logger.info('rtol: {}, atol: {}'.format(rtol, atol))
try:
for i in range(len(all_outputs)):
ref_outputs = all_ref_outputs[i]
outputs = all_outputs[i]
for i in range(len(all_outputs)):
ref_outputs = all_ref_outputs[i]
outputs = all_outputs[i]
for j in range(len(outputs)):
ref_output = ref_outputs[j]
output = outputs[j]
for j in range(len(outputs)):
ref_output = ref_outputs[j]
output = outputs[j]
# Compare the results with reference outputs
for ref_o, o in zip(ref_output, output):
# abs(desired-actual) < rtol * abs(desired) + atol
# Compare the results with reference outputs
for ref_o, o in zip(ref_output, output):
# abs(desired-actual) < rtol * abs(desired) + atol
try:
logger.info("Output shape{} input shape{}".format(ref_output.shape, output.shape))
np.testing.assert_allclose(ref_o, o, rtol, atol)
except Exception as e:
logger.error(e)
return False, e
except Exception as e:
if percentage_in_allowed_threshold(e, percent_mismatch):
continue
logger.error(e)
return False, e
logger.info('ONNX Runtime outputs are similar to reference outputs!')
return True, None
@ -482,7 +426,7 @@ def remove_profiling_files(path):
for f in files:
if "custom_test_data" in f:
continue
subprocess.Popen(["rm","-rf", f], stdout=subprocess.PIPE)
subprocess.Popen(["sudo","rm","-rf", f], stdout=subprocess.PIPE)
def update_fail_report(fail_results, model, ep, e_type, e):
@ -662,56 +606,59 @@ def write_map_to_file(result, file_name):
file.write(json.dumps(existed_result)) # use `json.loads` to do the reverse
def get_system_info():
info = {}
info["cuda"] = get_cuda_version()
info["trt"] = get_trt_version()
p = subprocess.Popen(["cat", "/etc/os-release"], stdout=subprocess.PIPE)
stdout, sterr = p.communicate()
stdout = stdout.decode("ascii").strip()
stdout = stdout.split("\n")[:2]
def get_cuda_version():
nvidia_strings = get_output(["nvidia-smi"])
version = re.search(r'CUDA Version: \d\d\.\d', nvidia_strings).group(0)
return version
def get_trt_version():
nvidia_strings = get_output(["dpkg", "-l"])
version = re.search(r'nvinfer.*\d\.\d\.\d\-\d', nvidia_strings).group(0)
return version
def get_linux_distro():
linux_strings = get_output(["cat", "/etc/os-release"])
stdout = linux_strings.split("\n")[:2]
infos = []
for row in stdout:
row = re.sub('=', ': ', row)
row = re.sub('"', '', row)
infos.append(row)
info["linux_distro"] = infos
return infos
p = subprocess.Popen(["lscpu"], stdout=subprocess.PIPE)
stdout, sterr = p.communicate()
stdout = stdout.decode("ascii").strip()
stdout = stdout.split("\n")
infos = []
for row in stdout:
if "mode" in row or "Arch" in row or "name" in row:
# row = row.replace(":\s+", ": ")
row = re.sub(': +', ': ', row)
infos.append(row)
info["cpu_info"] = infos
p1 = subprocess.Popen(["lspci", "-v"], stdout=subprocess.PIPE)
p2 = subprocess.Popen(["grep", "NVIDIA"], stdin=p1.stdout, stdout=subprocess.PIPE)
stdout, sterr = p2.communicate()
stdout = stdout.decode("ascii").strip()
stdout = stdout.split("\n")
infos = []
for row in stdout:
row = re.sub('.*:', '', row)
infos.append(row)
info["gpu_info"] = infos
p = subprocess.Popen(["cat", "/proc/meminfo"], stdout=subprocess.PIPE)
stdout, sterr = p.communicate()
stdout = stdout.decode("ascii").strip()
stdout = stdout.split("\n")
def get_memory_info():
mem_strings = get_output(["cat", "/proc/meminfo"])
stdout = mem_strings.split("\n")
infos = []
for row in stdout:
if "Mem" in row:
row = re.sub(': +', ': ', row)
infos.append(row)
info["memory"] = infos
return infos
def get_cpu_info():
cpu_strings = get_output(["lscpu"])
stdout = cpu_strings.split("\n")
infos = []
for row in stdout:
if "mode" in row or "Arch" in row or "name" in row:
row = re.sub(': +', ': ', row)
infos.append(row)
return infos
def get_gpu_info():
info = get_output(["lspci", "-v"])
infos = re.findall('NVIDIA.*', info)
return infos
def get_system_info():
info = {}
info["cuda"] = get_cuda_version()
info["trt"] = get_trt_version()
info["linux_distro"] = get_linux_distro()
info["cpu_info"] = get_cpu_info()
info["gpu_info"] = get_gpu_info()
info["memory"] = get_memory_info()
return info
@ -772,10 +719,6 @@ def parse_models_info_from_directory(path, models):
model_name = model_name + '_' + os.path.split(os.path.split(path)[0])[-1] # get opset version as model_name
model_path = find_model_path(path)
if not model_path:
logger.info("Can't find model in " + path)
return
model = {}
model["model_name"] = model_name
model["model_path"] = model_path
@ -923,10 +866,6 @@ def run_onnxruntime(args, models):
os.chdir(path)
path = os.getcwd()
# cleanup files before running a new inference
if args.running_mode == "validate":
remove_profiling_files(path)
inputs = []
ref_outputs = []
all_inputs_shape = [] # use for standalone trt
@ -1095,7 +1034,7 @@ def run_onnxruntime(args, models):
try:
ort_outputs = inference_ort_and_get_prediction(name, sess, inputs)
status = validate(ref_outputs, ort_outputs)
status = validate(ref_outputs, ort_outputs, args.rtol, args.atol, args.percent_mismatch)
if not status[0]:
update_fail_model_map(model_to_fail_ep, name, ep, 'result accuracy issue', status[1])
continue
@ -1372,6 +1311,8 @@ def output_latency(results, csv_filename):
row = [key,
cpu_average,
cpu_90_percentile,
cuda_average,
cuda_90_percentile,
trt_average,
@ -1515,6 +1456,11 @@ def parse_arguments():
parser.add_argument("--trtexec", required=False, default=None, help="trtexec executable path.")
# Validation options
parser.add_argument("--percent_mismatch", required=False, default=20.0, help="Allowed percentage of mismatched elements in validation.")
parser.add_argument("--rtol", required=False, default=0, help="Relative tolerance for validating outputs.")
parser.add_argument("--atol", required=False, default=20, help="Absolute tolerance for validating outputs.")
parser.add_argument("-t",
"--test_times",
required=False,

43
onnxruntime/python/tools/tensorrt/perf/benchmark_wrapper.py
View file

@ -6,7 +6,6 @@ import argparse
import copy
import json
import re
import sys
import pprint
from benchmark import *
from perf_utils import get_latest_commit_hash
@ -33,12 +32,11 @@ def main():
model_to_fail_ep = {}
commit = get_latest_commit_hash()
benchmark_fail_csv = 'fail_' + commit + '.csv'
benchmark_metrics_csv = 'metrics_' + commit + '.csv'
benchmark_success_csv = 'success_' + commit + '.csv'
benchmark_latency_csv = 'latency_' + commit + '.csv'
benchmark_status_csv = 'status_' + commit + '.csv'
benchmark_fail_csv = 'fail.csv'
benchmark_metrics_csv = 'metrics.csv'
benchmark_success_csv = 'success.csv'
benchmark_latency_csv = 'latency.csv'
benchmark_status_csv = 'status.csv'
for model, model_info in models.items():
logger.info("\n" + "="*40 + "="*len(model))
@ -49,8 +47,10 @@ def main():
model_list_file = os.path.join(os.getcwd(), model +'.json')
write_model_info_to_file([model_info], model_list_file)
ep_list = get_ep_list(args.comparison)
if args.ep:
ep_list = [args.ep]
else:
ep_list = get_ep_list(args.comparison)
for ep in ep_list:
if args.running_mode == "validate":
p = subprocess.run(["python3",
@ -91,24 +91,25 @@ def main():
Path(path).mkdir(parents=True, exist_ok=True)
if args.running_mode == "validate":
logger.info("\n==========================================================")
logger.info("\n=========================================================")
logger.info("========== Failing Models/EPs (accumulated) ==============")
logger.info("==========================================================")
if os.path.exists(FAIL_MODEL_FILE) or len(model_to_fail_ep) > 1:
model_to_fail_ep = read_map_from_file(FAIL_MODEL_FILE)
output_fail(model_to_fail_ep, os.path.join(path, benchmark_fail_csv))
logger.info("\nSaved model fail results to {}".format(benchmark_fail_csv))
logger.info(model_to_fail_ep)
logger.info("\n=========================================")
logger.info("========== Models/EPs metrics ==========")
logger.info("=========== Models/EPs metrics ==========")
logger.info("=========================================")
if os.path.exists(METRICS_FILE):
model_to_metrics = read_map_from_file(METRICS_FILE)
output_metrics(model_to_metrics, os.path.join(path, benchmark_metrics_csv))
logger.info("\nSaved model metrics results to {}".format(benchmark_metrics_csv))
elif args.running_mode == "benchmark":
logger.info("\n=======================================================")
logger.info("=========== Models/EPs Status (accumulated) ===========")
@ -124,28 +125,30 @@ def main():
model_fail = read_map_from_file(FAIL_MODEL_FILE)
is_fail = True
model_status = build_status(model_status, model_fail, is_fail)
pretty_print(pp, model_status)
pp.pprint(model_status)
output_status(model_status, os.path.join(path, benchmark_status_csv))
logger.info("\nSaved model status results to {}".format(benchmark_status_csv))
logger.info("\n=======================================================")
logger.info("\n=========================================================")
logger.info("=========== Models/EPs latency (accumulated) ===========")
logger.info("=======================================================")
logger.info("=========================================================")
if os.path.exists(LATENCY_FILE):
model_to_latency = read_map_from_file(LATENCY_FILE)
add_improvement_information(model_to_latency)
pretty_print(pp, model_to_latency)
output_latency(model_to_latency, os.path.join(path, benchmark_latency_csv))
pp.pprint(model_to_latency)
logger.info("\nSaved model status results to {}".format(benchmark_latency_csv))
logger.info("\n===========================================")
logger.info("=========== System information ===========")
logger.info("===========================================")
info = get_system_info()
pp.pprint(info)
pretty_print(pp, info)
if __name__ == "__main__":
main()

4
onnxruntime/python/tools/tensorrt/perf/build/Dockerfile.tensorrt-perf Normal file
View file

@ -0,0 +1,4 @@
FROM onnxruntime-trt
RUN apt-get update &&\
apt-get -y install libprotobuf-dev protobuf-compiler pciutils &&\
pip install coloredlogs numpy flake8 onnx Cython onnxmltools sympy packaging

14
onnxruntime/python/tools/tensorrt/perf/build/build_images.sh Executable file
View file

@ -0,0 +1,14 @@
#!/bin/bash
while getopts o:p:b:i: parameter
do case "${parameter}"
in
o) ORT_DOCKERFILE_PATH=${OPTARG};;
p) PERF_DOCKERFILE_PATH=${OPTARG};;
b) ORT_BRANCH=${OPTARG};;
i) IMAGE_NAME=${OPTARG};;
esac
done
sudo docker build --no-cache -t onnxruntime-trt --build-arg ONNXRUNTIME_BRANCH=$ORT_BRANCH -f $ORT_DOCKERFILE_PATH ..
sudo docker build --no-cache -t $IMAGE_NAME -f $PERF_DOCKERFILE_PATH ..

46
onnxruntime/python/tools/tensorrt/perf/comparison_scripts/compare_latency.py Normal file
View file

@ -0,0 +1,46 @@
import pandas as pd
import numpy as np
import argparse
ep_map = {"cpu": "CPU", "cuda":"CUDA","trt": "TRT EP","native": "Standalone TRT"}
def parse_arguments():
# create parser
parser = argparse.ArgumentParser()
parser.add_argument("-p", "--prev", required=True, help="previous csv")
parser.add_argument("-c", "--current", required=True, help="current csv")
parser.add_argument("-o", "--output_csv", required=True, help="output different csv")
parser.add_argument("--ep", required=False, default="trt", choices=["cpu", "cuda", "trt", "native"], help="ep to capture regressions on")
parser.add_argument("--tolerance", required=False, default=0, help="allowed tolerance for latency comparison")
args = parser.parse_args()
return args
def get_table_condition(table, fp, ep, tol):
ep = ep_map[ep]
col1 = ep + " " + fp + " \nmean (ms)_x"
col2 = ep + " " + fp + " \nmean (ms)_y"
condition = table[col1] > (table[col2] + tol)
return condition
def main():
args = parse_arguments()
a = pd.read_csv(args.prev)
b = pd.read_csv(args.current)
common = a.merge(b, on=['Model'])
condition_fp32 = get_table_condition(common, "fp32", args.ep, args.tolerance)
condition_fp16 = get_table_condition(common, "fp16", args.ep, args.tolerance)
common['greater'] = np.where((condition_fp32 | condition_fp16), True, False)
greater = common[common['greater'] == True].drop(['greater'], axis=1)
# arrange columns
keys = list(greater.keys().sort_values())
keys.insert(0, keys.pop(keys.index('Model')))
greater = greater[keys]
greater.to_csv(args.output_csv)
if __name__=='__main__':
main()

22
onnxruntime/python/tools/tensorrt/perf/comparison_scripts/new_failures.py Normal file
View file

@ -0,0 +1,22 @@
import pandas as pd
import argparse
def parse_arguments():
# create parser
parser = argparse.ArgumentParser()
parser.add_argument("-p", "--prev", required=True, help="previous csv")
parser.add_argument("-c", "--current", required=True, help="current csv")
parser.add_argument("-o", "--output_csv", required=True, help="output different csv")
args = parser.parse_args()
return args
def main():
args = parse_arguments()
a = pd.read_csv(args.prev)
b = pd.read_csv(args.current)
common = b.merge(a, on=['model','ep','error type','error message'])
diff = b.append(common, ignore_index=True).drop_duplicates(['model', 'ep', 'error type', 'error message'], keep=False).loc[:b.index.max()]
diff.to_csv(args.output_csv)
if __name__=='__main__':
main()

14
onnxruntime/python/tools/tensorrt/perf/ort_build_latest.py
View file

@ -1,6 +1,7 @@
import os
import subprocess
import argparse
import tarfile
from perf_utils import get_latest_commit_hash
def parse_arguments():
@ -34,12 +35,13 @@ def install_new_ort_wheel(ort_master_path):
def main():
args = parse_arguments()
p1 = subprocess.Popen(["sudo", "wget", "https://cmake.org/files/v3.17/cmake-3.17.4-Linux-x86_64.tar.gz"])
p1.wait()
p1 = subprocess.Popen(["tar", "zxvf", "cmake-3.17.4-Linux-x86_64.tar.gz"])
p1.wait()
cmake_tar = "cmake-3.17.4-Linux-x86_64.tar.gz"
if not os.path.exists(cmake_tar):
p = subprocess.run(["wget", "-c", "https://cmake.org/files/v3.17/" + cmake_tar], check=True)
tar = tarfile.open(cmake_tar)
tar.extractall()
tar.close()
os.environ["PATH"] = os.path.join(os.path.abspath("cmake-3.17.4-Linux-x86_64"), "bin") + ":" + os.environ["PATH"]
os.environ["CUDACXX"] = os.path.join(args.cuda_home, "bin", "nvcc")

58
onnxruntime/python/tools/tensorrt/perf/perf.sh
View file

@ -1,9 +1,21 @@
#!/bin/bash
while getopts d:o:m: parameter
do case "${parameter}"
in
d) PERF_DIR=${OPTARG};;
o) OPTION=${OPTARG};;
m) MODEL_PATH=${OPTARG};;
esac
done
# metadata
FAIL_MODEL_FILE=".fail_model_map"
LATENCY_FILE=".latency_map"
METRICS_FILE=".metrics_map"
PROFILE="*onnxruntime_profile*"
# files to download info
SYMBOLIC_SHAPE_INFER="symbolic_shape_infer.py"
@ -17,49 +29,21 @@ cleanup_files() {
rm -f $METRICS_FILE
rm -f $SYMBOLIC_SHAPE_INFER
rm -f $FLOAT_16
rm -rf result/$OPTION
find -name $PROFILE -delete
}
download_files() {
sudo wget -c $SYMBOLIC_SHAPE_INFER_LINK
sudo wget -c $FLOAT_16_LINK
wget --no-check-certificate -c $SYMBOLIC_SHAPE_INFER_LINK
wget --no-check-certificate -c $FLOAT_16_LINK
}
update_files() {
setup() {
cd $PERF_DIR
cleanup_files
download_files
}
# many models
if [ "$1" == "many-models" ]
then
update_files
python3 benchmark_wrapper.py -r validate -m /home/hcsuser/mount/many-models -o result/"$1"
python3 benchmark_wrapper.py -r benchmark -i random -t 10 -m /home/hcsuser/mount/many-models -o result/"$1"
fi
# ONNX model zoo
if [ "$1" == "onnx-zoo-models" ]
then
MODEL_LIST="model_list.json"
update_files
python3 benchmark_wrapper.py -r validate -m $MODEL_LIST -o result/"$1"
python3 benchmark_wrapper.py -r benchmark -i random -t 10 -m $MODEL_LIST -o result/"$1"
fi
# 1P models
if [ "$1" == "partner-models" ]
then
MODEL_LIST="partner_model_list.json"
update_files
python3 benchmark_wrapper.py -r validate -m $MODEL_LIST -o result/"$1"
python3 benchmark_wrapper.py -r benchmark -i random -t 10 -m $MODEL_LIST -o result/"$1"
fi
# Test models
if [ "$1" == "selected-models" ]
then
MODEL_LIST="selected_models.json"
update_files
python3 benchmark_wrapper.py -r validate -m $MODEL_LIST -o result/"$1"
python3 benchmark_wrapper.py -r benchmark -i random -t 1 -m $MODEL_LIST -o result/"$1"
fi
setup
python3 benchmark_wrapper.py -r validate -m $MODEL_PATH -o result/$OPTION
python3 benchmark_wrapper.py -r benchmark -i random -t 10 -m $MODEL_PATH -o result/$OPTION

14
onnxruntime/python/tools/tensorrt/perf/perf_utils.py
View file

@ -4,20 +4,28 @@ import pprint
import logging
import coloredlogs
import re
import sys
debug = False
debug_verbose = False
def get_output(command):
p = subprocess.run(command, check=True, stdout=subprocess.PIPE)
output = p.stdout.decode("ascii").strip()
return output
def find(regex_string):
import glob
results = glob.glob(regex_string)
results.sort()
return results
def pretty_print(pp, json_object):
pp.pprint(json_object)
sys.stdout.flush()
def get_latest_commit_hash():
p1 = subprocess.Popen(["git", "rev-parse", "--short", "HEAD"], stdout = subprocess.PIPE)
stdout, sterr = p1.communicate()
commit = stdout.decode("utf-8").strip()
commit = get_output(["git", "rev-parse", "--short", "HEAD"])
return commit
def parse_single_file(f):

36
onnxruntime/python/tools/tensorrt/perf/run_perf_docker.sh Executable file
View file

@ -0,0 +1,36 @@
#!/bin/bash
# Parse Arguments
while getopts d:o:m: parameter
do case "${parameter}"
in
d) DOCKER_IMAGE=${OPTARG};;
o) OPTION=${OPTARG};;
m) MODEL_PATH=${OPTARG};;
esac
done
# Variables
MACHINE_PERF_DIR=/home/hcsuser/perf/
DOCKER_PERF_DIR=/usr/share/perf/
PERF_SCRIPT=$DOCKER_PERF_DIR'perf.sh'
VOLUME=$MACHINE_PERF_DIR:$DOCKER_PERF_DIR
# Add Remaining Variables
if [ $OPTION == "onnx-zoo-models" ]
then
MODEL_PATH=model.json
fi
if [ $OPTION == "many-models" ]
then
MODEL_PATH=/usr/share/mount/many-models
VOLUME=$VOLUME' -v /home/hcsuser/mount/test:/usr/share/mount/many-models'
fi
if [ $OPTION == "partner-models" ]
then
MODEL_PATH=partner_model_list.json
fi
sudo docker run --gpus all -v $VOLUME $DOCKER_IMAGE /bin/bash $PERF_SCRIPT -d $DOCKER_PERF_DIR -o $OPTION -m $MODEL_PATH

31
onnxruntime/python/tools/tensorrt/perf/run_perf_machine.sh Executable file
View file

@ -0,0 +1,31 @@
#!/bin/bash
# Parse Arguments
while getopts d:o:m: parameter
do case "${parameter}"
in
o) OPTION=${OPTARG};;
m) MODEL_PATH=${OPTARG};;
esac
done
# Variables
PERF_DIR=/home/hcsuser/perf/
# Select models to be tested or run selected-models
if [ $OPTION == "onnx-zoo-models" ]
then
MODEL_PATH='model.json'
fi
if [ $OPTION == "many-models" ]
then
MODEL_PATH=/usr/share/mount/many-models
fi
if [ $OPTION == "partner-models" ]
then
MODEL_PATH='partner_model_list.json'
fi
./perf.sh -d $PERF_DIR -o $OPTION -m $MODEL_PATH

5
onnxruntime/python/tools/tensorrt/perf/setup_scripts/setup_many_models.sh Executable file
View file

@ -0,0 +1,5 @@
# !/bin/bash
while true;do wget -T 15 -c "$1" && break;done
FILE="$(basename -- "$1")"
unzip $FILE.zip

71
onnxruntime/python/tools/tensorrt/perf/setup_scripts/setup_onnx_zoo.py Normal file
View file

@ -0,0 +1,71 @@
import os
import wget
import tarfile
import json
def get_tar_file(link):
file_name = link.split("/")[-1]
return file_name
def create_model_folder(model):
os.mkdir(model)
def extract_and_get_files(file_name):
model_folder = file_name.replace(".tar.gz", "") + '/'
create_model_folder(model_folder)
model_tar = tarfile.open(file_name)
model_tar.extractall(model_folder)
file_list = model_tar.getnames()
file_list.sort()
model_tar.close()
return model_folder, file_list
def download_model(link):
file_name = get_tar_file(link)
wget.download(link)
model_folder, file_list = extract_and_get_files(file_name)
return model_folder, file_list
def get_model_path(file_list):
for file_name in file_list:
if ".onnx" in file_name:
return file_name
def get_test_path(model_path):
model_filename = os.path.basename(model_path)
test_path = model_path.split(model_filename)[0]
return test_path
def create_model_object(model, folder, model_file_path, test_path):
model_dict = {}
model_dict["model_name"] = model
model_dict["working_directory"] = "./models/" + folder
model_dict["model_path"] = "./" + model_file_path
model_dict["test_data_path"] = "./" + test_path
return model_dict
def get_model_info(link):
model_folder, file_list = download_model(link)
model = model_folder[:-1]
model_file_path = get_model_path(file_list)
test_path = get_test_path(model_file_path)
model_info = create_model_object(model, model_folder, model_file_path, test_path)
return model_info
def write_json(models):
model_json = json.dumps(models, indent=4)
with open('model_list.json', 'w') as fp:
fp.write(models_json)
def main():
links = []
with open('links.txt', 'r') as fh:
links = [link.rstrip() for link in fh.readlines()]
model_list = []
for link in links:
model_list.append(get_model_info(link))
write_json(model_list)
if __name__ == "__main__":
main()

4
onnxruntime/python/tools/transformers/README.md
View file

@ -210,7 +210,7 @@ For GPU, please append --use_gpu to the command.
bert_perf_test.py can be used to check the BERT model inference performance. Below are examples:
```console
python -m onnxruntime.transformers.bert_perf_test --model optimized_model_cpu.onnx --batch_size 1 --sequence_length 128 --samples 100 --test_times 10 --inclusive
python -m onnxruntime.transformers.bert_perf_test --model optimized_model_cpu.onnx --batch_size 1 --sequence_length 128
```
For GPU, please append --use_gpu to the command.
@ -219,7 +219,7 @@ After test is finished, a file like perf_results_CPU_B1_S128_<date_time>.txt or
## Profiling
profiler.py can be used to run profiling on a transformer model. It can help figure out the bottleneck of a model, and time spent on a node or subgraph.
profiler.py can be used to run profiling on a transformer model. It can help figure out the bottleneck of a model, and CPU time spent on a node or subgraph.
Examples commands:

3
onnxruntime/python/tools/transformers/benchmark.py
View file

@ -80,9 +80,6 @@ def run_onnxruntime(use_gpu, model_names, model_class, precision, num_threads, b
)
return results
if (not use_gpu) and ('CUDAExecutionProvider' in onnxruntime.get_available_providers()):
logger.warning("Please install onnxruntime package instead of onnxruntime-gpu to get best cpu performance.")
for model_name in model_names:
all_input_names = MODELS[model_name][0]
for num_inputs in input_counts:

7
onnxruntime/python/tools/transformers/benchmark_gpt2.py
View file

@ -16,6 +16,7 @@ import argparse
import logging
import torch
import onnx
from packaging import version
from transformers import AutoConfig
from gpt2_helper import Gpt2Helper, MODEL_CLASSES, DEFAULT_TOLERANCE, PRETRAINED_GPT2_MODELS
from quantize_helper import QuantizeHelper
@ -113,6 +114,10 @@ def parse_arguments(argv=None):
def main(args):
from transformers import __version__ as transformers_version
if version.parse(transformers_version) < version.parse("3.1.0"): # past_key_values name does not exist in 3.0.2 or older
raise RuntimeError("This tool requires transformers 3.1.0 or later.")
logger.info(f"Arguments:{args}")
if args.precision == Precision.FLOAT16:
assert args.optimize_onnx and args.use_gpu, "fp16 requires --optimize_onnx --use_gpu"
@ -279,7 +284,7 @@ def main(args):
return csv_filename
if __name__ == '__main__':
if __name__ == '__main__':
args = parse_arguments()
setup_logger(args.verbose)
main(args)

87
onnxruntime/python/tools/transformers/benchmark_helper.py
View file

@ -235,3 +235,90 @@ def allocateOutputBuffers(output_buffers, output_buffer_max_sizes, device):
for i in output_buffer_max_sizes:
output_buffers.append(torch.empty(i, dtype=torch.float32, device=device))
def set_random_seed(seed=123):
"""Set random seed manully to get deterministic results"""
import random
random.seed(seed)
numpy.random.seed(seed)
torch.manual_seed(seed)
torch.cuda.manual_seed(seed)
torch.cuda.manual_seed_all(seed)
#torch.backends.cudnn.enabled = False
#torch.backends.cudnn.benchmark = False
#torch.backends.cudnn.deterministic = True
def measure_memory(is_gpu, func):
import os
import psutil
from time import sleep
class MemoryMonitor:
def __init__(self, keep_measuring=True):
self.keep_measuring = keep_measuring
def measure_cpu_usage(self):
max_usage = 0
while True:
max_usage = max(max_usage, psutil.Process(os.getpid()).memory_info().rss / 1024**2)
sleep(0.005) # 5ms
if not self.keep_measuring:
break
return max_usage
def measure_gpu_usage(self):
from py3nvml.py3nvml import nvmlInit, nvmlDeviceGetCount, nvmlDeviceGetHandleByIndex, \
nvmlDeviceGetMemoryInfo, nvmlDeviceGetName, nvmlShutdown, NVMLError
max_gpu_usage = []
gpu_name = []
try:
nvmlInit()
deviceCount = nvmlDeviceGetCount()
max_gpu_usage = [0 for i in range(deviceCount)]
gpu_name = [nvmlDeviceGetName(nvmlDeviceGetHandleByIndex(i)) for i in range(deviceCount)]
while True:
for i in range(deviceCount):
info = nvmlDeviceGetMemoryInfo(nvmlDeviceGetHandleByIndex(i))
max_gpu_usage[i] = max(max_gpu_usage[i], info.used / 1024**2)
sleep(0.005) # 5ms
if not self.keep_measuring:
break
nvmlShutdown()
return [{
"device_id": i,
"name": gpu_name[i],
"max_used_MB": max_gpu_usage[i]
} for i in range(deviceCount)]
except NVMLError as error:
if not self.silent:
self.logger.error("Error fetching GPU information using nvml: %s", error)
return None
monitor = MemoryMonitor(False)
memory_before_test = monitor.measure_gpu_usage() if is_gpu else monitor.measure_cpu_usage()
from concurrent.futures import ThreadPoolExecutor
with ThreadPoolExecutor() as executor:
monitor = MemoryMonitor()
mem_thread = executor.submit(monitor.measure_gpu_usage if is_gpu else monitor.measure_cpu_usage)
try:
fn_thread = executor.submit(func)
result = fn_thread.result()
finally:
monitor.keep_measuring = False
max_usage = mem_thread.result()
if is_gpu:
print(f"GPU memory usage: before={memory_before_test} peak={max_usage}")
if len(memory_before_test) >= 1 and len(max_usage) >= 1:
before = memory_before_test[0]["max_used_MB"]
after = max_usage[0]["max_used_MB"]
return after - before
else:
return None
else:
print(f"CPU memory usage: before={memory_before_test:.1f} MB, peak={max_usage:.1f} MB")
return max_usage - memory_before_test

241
onnxruntime/python/tools/transformers/bert_perf_test.py
View file

@ -35,44 +35,10 @@ class TestSetting:
sequence_length: int
test_cases: int
test_times: int
contiguous: bool
use_gpu: bool
warmup: bool
omp_num_threads: int
omp_wait_policy: str
intra_op_num_threads: int
seed: int
verbose: bool
contiguous: bool
inclusive: bool
extra_latency: float = 0
def get_setting(self) -> str:
return f"batch_size={self.batch_size},sequence_length={self.sequence_length},test_cases={self.test_cases},test_times={self.test_times},contiguous={self.contiguous},use_gpu={self.use_gpu},warmup={self.warmup}"
def check(self, intra_op_threads, omp_threads, omp_policy) -> bool:
if intra_op_threads is None:
if self.intra_op_num_threads is not None and self.intra_op_num_threads > 0:
return False
else:
assert intra_op_threads > 0
if not (self.intra_op_num_threads is None or self.intra_op_num_threads == intra_op_threads):
return False
if omp_threads is None:
if self.omp_num_threads is not None and self.omp_num_threads > 0:
return False
else:
assert omp_threads > 0
if not (self.omp_num_threads is None or self.omp_num_threads == omp_threads):
return False
if self.omp_wait_policy is not None:
if omp_policy != self.omp_wait_policy:
return False
return True
@dataclass
class ModelSetting:
@ -84,22 +50,17 @@ class ModelSetting:
def create_session(model_path, use_gpu, intra_op_num_threads, graph_optimization_level=None):
# Import onnxruntime shall be after OpenMP environment variable setting.
# So we put the import in function to delay importing instead of top of this script.
import onnxruntime
if use_gpu and ('CUDAExecutionProvider' not in onnxruntime.get_available_providers()):
print(
"Warning: Please install onnxruntime-gpu package instead of onnxruntime, and use a machine with GPU for testing gpu performance."
)
elif (not use_gpu) and ('CUDAExecutionProvider' in onnxruntime.get_available_providers()):
print("Warning: Please install onnxruntime package instead of onnxruntime-gpu to get best cpu performance.")
if intra_op_num_threads is None and graph_optimization_level is None:
session = onnxruntime.InferenceSession(model_path)
else:
execution_providers = ['CPUExecutionProvider'
] if not use_gpu else ['CUDAExecutionProvider', 'CPUExecutionProvider']
execution_providers = ['CPUExecutionProvider'] if not use_gpu else ['CUDAExecutionProvider', 'CPUExecutionProvider']
sess_options = onnxruntime.SessionOptions()
sess_options.execution_mode = onnxruntime.ExecutionMode.ORT_SEQUENTIAL
@ -127,8 +88,8 @@ def create_session(model_path, use_gpu, intra_op_num_threads, graph_optimization
return session
def onnxruntime_inference(session, all_inputs, output_names, warmup=True):
if warmup and len(all_inputs) > 0:
def onnxruntime_inference(session, all_inputs, output_names):
if len(all_inputs) > 0:
# Use a random input as warm up.
session.run(output_names, random.choice(all_inputs))
@ -142,57 +103,16 @@ def onnxruntime_inference(session, all_inputs, output_names, warmup=True):
latency_list.append(latency)
return results, latency_list
def get_contiguous_inputs(all_inputs):
"""
Convert input to be contiguous.
"""
contiguous_inputs = []
start_time = timeit.default_timer()
for test_case_id, inputs in enumerate(all_inputs):
real_inputs = {}
for key, value in inputs.items():
real_inputs[key] = np.ascontiguousarray(value)
contiguous_inputs.append(real_inputs)
latency = timeit.default_timer() - start_time
average_latency_ms = latency / len(contiguous_inputs) * 1000
return contiguous_inputs, average_latency_ms
def to_string(model_path, session, test_setting):
sess_options = session.get_session_options()
option = "model={}".format(os.path.basename(model_path))
option += ",graph_optimization_level={},intra_op_num_threads={}".format(sess_options.graph_optimization_level,
option = "model={},".format(os.path.basename(model_path))
option += "graph_optimization_level={},intra_op_num_threads={},".format(sess_options.graph_optimization_level,
sess_options.intra_op_num_threads).replace(
'GraphOptimizationLevel.ORT_', '')
option += ",OMP_NUM_THREADS={}".format(os.environ["OMP_NUM_THREADS"] if "OMP_NUM_THREADS" in os.environ else "")
option += ",OMP_WAIT_POLICY={}".format(os.environ["OMP_WAIT_POLICY"] if "OMP_WAIT_POLICY" in os.environ else "")
option += ",{}".format(test_setting.get_setting())
option += f"batch_size={test_setting.batch_size},sequence_length={test_setting.sequence_length},test_cases={test_setting.test_cases},test_times={test_setting.test_times},use_gpu={test_setting.use_gpu}"
return option
def setup_openmp_environ(omp_num_threads, omp_wait_policy):
if omp_num_threads is None:
if "OMP_NUM_THREADS" in os.environ:
del os.environ["OMP_NUM_THREADS"]
else:
os.environ["OMP_NUM_THREADS"] = str(omp_num_threads)
if omp_wait_policy is None:
if "OMP_WAIT_POLICY" in os.environ:
del os.environ["OMP_WAIT_POLICY"]
else:
assert omp_wait_policy in ["ACTIVE", "PASSIVE"], f"{omp_wait_policy} is not a valid policy"
os.environ["OMP_WAIT_POLICY"] = omp_wait_policy
def run_one_test(model_setting, test_setting, perf_results, all_inputs, intra_op_num_threads, omp_num_threads,
omp_wait_policy):
# Environment variable shall be set before import onnxruntime.
setup_openmp_environ(omp_num_threads, omp_wait_policy)
def run_one_test(model_setting, test_setting, perf_results, all_inputs, intra_op_num_threads):
session = create_session(model_setting.model_path, test_setting.use_gpu, intra_op_num_threads,
model_setting.opt_level)
output_names = [output.name for output in session.get_outputs()]
@ -206,11 +126,11 @@ def run_one_test(model_setting, test_setting, perf_results, all_inputs, intra_op
all_latency_list = []
for i in range(test_setting.test_times):
results, latency_list = onnxruntime_inference(session, all_inputs, output_names, test_setting.warmup)
results, latency_list = onnxruntime_inference(session, all_inputs, output_names)
all_latency_list.extend(latency_list)
# latency in miliseconds
latency_ms = np.array(all_latency_list) * 1000 + test_setting.extra_latency
latency_ms = np.array(all_latency_list) * 1000
average_latency = statistics.mean(latency_ms)
latency_50 = np.percentile(latency_ms, 50)
@ -226,91 +146,31 @@ def run_one_test(model_setting, test_setting, perf_results, all_inputs, intra_op
format(throughput, '.2f')))
def launch_test(model_setting, test_setting, perf_results, all_inputs, intra_op_num_threads, omp_num_threads,
omp_wait_policy):
if not test_setting.check(intra_op_num_threads, omp_num_threads, omp_wait_policy):
return
def launch_test(model_setting, test_setting, perf_results, all_inputs, intra_op_num_threads):
process = multiprocessing.Process(target=run_one_test,
args=(model_setting, test_setting, perf_results, all_inputs, intra_op_num_threads,
omp_num_threads, omp_wait_policy))
args=(model_setting, test_setting, perf_results, all_inputs, intra_op_num_threads))
process.start()
process.join()
def run_perf_tests(model_setting, test_setting, perf_results, test_all, all_inputs):
def run_perf_tests(model_setting, test_setting, perf_results, all_inputs):
if (test_setting.intra_op_num_threads is not None):
launch_test(model_setting, test_setting, perf_results, all_inputs, test_setting.intra_op_num_threads)
return
cpu_count = psutil.cpu_count(logical=False)
logical_cores = psutil.cpu_count(logical=True)
candidate_threads = list(set([1, logical_cores, cpu_count]))
if (test_setting.intra_op_num_threads is not None) or (test_setting.omp_num_threads is not None):
if test_setting.intra_op_num_threads is not None:
intra_op_threads = [test_setting.intra_op_num_threads]
else:
intra_op_threads = [None] + candidate_threads
if test_setting.omp_num_threads is not None:
omp_threads = [test_setting.omp_num_threads]
else:
omp_threads = [None] + candidate_threads
if test_setting.omp_wait_policy is not None:
omp_policies = [test_setting.omp_wait_policy]
else:
omp_policies = [None, 'PASSIVE', 'ACTIVE']
for it in intra_op_threads:
for ot in omp_threads:
for op in omp_policies:
launch_test(model_setting, test_setting, perf_results, all_inputs, it, ot, op)
return
# Test a setting without any setting as baseline 1.
launch_test(model_setting, test_setting, perf_results, all_inputs, None, None, None)
if not test_setting.use_gpu:
# For CPU: intra_op_num_threads = 1, omp_num_threads=None, omp_wait_policy=None
# Another setting without environment variable as baseline 2.
launch_test(model_setting, test_setting, perf_results, all_inputs, 1, None, None)
else:
# For GPU, we test two more settings by default:
# (1) intra_op_num_threads = 1, omp_num_threads=cpu_count, omp_wait_policy=PASSIVE
# (2) intra_op_num_threads = logical_cores, omp_num_threads=1, omp_wait_policy=ACTIVE
launch_test(model_setting, test_setting, perf_results, all_inputs, 1, cpu_count, 'PASSIVE')
launch_test(model_setting, test_setting, perf_results, all_inputs, logical_cores, 1, 'ACTIVE')
# GPU latency is not sensitive to these settings. No need to test many combinations.
# Skip remaining settings for GPU without --all flag.
if test_setting.use_gpu and not test_all:
return
candidate_threads = list(set([logical_cores, cpu_count]))
for i in range(1, min(16, logical_cores)):
if i not in candidate_threads:
candidate_threads.append(i)
candidate_threads.sort(reverse=True)
for intra_op_num_threads in candidate_threads:
for omp_num_threads in candidate_threads:
# skip settings that are very slow
if intra_op_num_threads == 1 and omp_num_threads == 1 and logical_cores != 1:
continue
# When logical and physical cores are not the same, there are many combinations.
# Remove some settings are not good normally.
if logical_cores > cpu_count:
if omp_num_threads == logical_cores and intra_op_num_threads != 1:
continue
if intra_op_num_threads == logical_cores and omp_num_threads != 1:
continue
if not test_all:
if intra_op_num_threads != 1 and omp_num_threads != 1:
continue
for omp_wait_policy in ['ACTIVE', 'PASSIVE']:
launch_test(model_setting, test_setting, perf_results, all_inputs, intra_op_num_threads,
omp_num_threads, omp_wait_policy)
def run_performance(model_setting, test_setting, perf_results, test_all):
launch_test(model_setting, test_setting, perf_results, all_inputs, intra_op_num_threads)
def run_performance(model_setting, test_setting, perf_results):
input_ids, segment_ids, input_mask = get_bert_inputs(model_setting.model_path, model_setting.input_ids_name,
model_setting.segment_ids_name, model_setting.input_mask_name)
@ -327,29 +187,25 @@ def run_performance(model_setting, test_setting, perf_results, test_all):
segment_ids,
input_mask,
random_mask_length=False)
if test_setting.contiguous:
all_inputs, contiguous_latency = get_contiguous_inputs(all_inputs)
print("Extra latency for converting inputs to contiguous: {} ms".format(format(contiguous_latency, '.2f')))
test_setting.extra_latency = contiguous_latency if test_setting.inclusive else 0
run_perf_tests(model_setting, test_setting, perf_results, test_all, all_inputs)
run_perf_tests(model_setting, test_setting, perf_results, all_inputs)
def parse_arguments():
parser = argparse.ArgumentParser()
parser.add_argument('--model', required=True, type=str, help="bert onnx model path")
parser.add_argument('--batch_size',
parser.add_argument('-b', '--batch_size',
required=True,
type=int,
nargs="+",
help="batch size of input. Allow one or multiple values in the range of [1, 128].")
parser.add_argument('--sequence_length', required=True, type=int, help="maximum sequence length of input")
parser.add_argument('-s', '--sequence_length', required=True, type=int, help="maximum sequence length of input")
parser.add_argument('--samples', required=False, type=int, default=10, help="number of samples to be generated")
parser.add_argument('--test_times',
parser.add_argument('-t', '--test_times',
required=False,
type=int,
default=0,
@ -375,40 +231,12 @@ def parse_arguments():
parser.add_argument('--use_gpu', required=False, action='store_true', help="use GPU")
parser.set_defaults(use_gpu=False)
parser.add_argument('--inclusive',
required=False,
action='store_true',
help="include the latency of converting array to contiguous")
parser.set_defaults(inclusive=False)
parser.add_argument('--all', required=False, action='store_true', help="test all candidate settings")
parser.set_defaults(all=False)
parser.add_argument('--omp_num_threads',
required=False,
type=int,
default=None,
help=">0, set OMP_NUM_THREADS value. 0, do not set")
parser.add_argument('--intra_op_num_threads',
parser.add_argument('-n', '--intra_op_num_threads',
required=False,
type=int,
default=None,
help=">=0, set intra_op_num_threads")
parser.add_argument('--omp_wait_policy',
required=False,
type=str,
default=None,
choices=['ACTIVE', 'PASSIVE'],
help="OMP_WAIT_POLICY")
parser.add_argument('--contiguous', required=False, action='store_true', help="contiguous input")
parser.set_defaults(contiguous=False)
parser.add_argument('--no_warmup', required=False, action='store_true', help="do not use one sample for warm-up.")
parser.set_defaults(no_warmup=False)
parser.add_argument('--input_ids_name', required=False, type=str, default=None, help="input name for input ids")
parser.add_argument('--segment_ids_name', required=False, type=str, default=None, help="input name for segment ids")
parser.add_argument('--input_mask_name',
@ -443,18 +271,13 @@ def main():
args.sequence_length,
args.samples,
args.test_times,
None, #contiguous
args.use_gpu,
not args.no_warmup,
args.omp_num_threads,
args.omp_wait_policy,
args.intra_op_num_threads,
args.seed,
args.verbose,
args.contiguous,
args.inclusive)
args.verbose)
print("test setting", test_setting)
run_performance(model_setting, test_setting, perf_results, args.all)
run_performance(model_setting, test_setting, perf_results)
# Sort the results so that the first one has smallest latency.
sorted_results = sorted(perf_results.items(), reverse=False, key=lambda x: x[1])

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