cherry pick outstanding commits (#7871)

* Fix bug in Transpose CUDA kernel (#7329)

* Fix permission error for ORTModule lock file (#7814)

* fix topo sort in quant tool (#7833)

* fix topo sort in quant tool

* add unit test and make the topo sort stable

* Relax tol for Conv1D fp16 test (#7844)

* Relax tol for Conv1D fp16 test

Co-authored-by: Sherlock Huang <bahuang@OrtTrainingDev3.af05slrtruoetgaxwwjv5nsq5e.px.internal.cloudapp.net>

* Resolve issue with wrapped ORTModule load_state_dict (#7847)

* Encapsulate children modules inside a ModuleAccessor object to prevent erroneuos iteration over children while loading the state dictionary

* Add named_models, models, apply methods, change ModuleAccessor to ModuleMetadata and modify unit tests

* Change ModuleMetadata module getter logic, raise NotImplementedError for add_modules

* Add comment explaining why overriding _load_from_state_dict method is needed

* fixed bugs in packed mode and enable pack mode tests in ci (#7848)

* fixed bugs in packed mode and enable pack mode tests in ci

* removed unnecessary space

* pr comments

* pr comments

* disable an average pool test

* try disabling another avg pool

* disable more avg pool tests

* disable maxpool tests

* add environment variable to control default training package's local version (#7849)

* [js] update documents (#7852)

* [js] update documents

* escape double quotes

* update operators.md

* resolve comments

* Support bool type for Pad CPU (#7856)

* Initial commit

* update

* nit

* Include ORT C/C++ API headers in the ORT Mobile AAR package (#7858)

* Add header files of ort c/c++ api to aar package

* Move header file selection to cmake based on EP choice

* fix duplicated node name (#7865)

* Clean up CPU kernel definition for opset 13 Pad (#7867)

Co-authored-by: Hariharan Seshadri <shariharan91@gmail.com>
Co-authored-by: Thiago Crepaldi <thiago.crepaldi@microsoft.com>
Co-authored-by: Yufeng Li <liyufeng1987@gmail.com>
Co-authored-by: Sherlock <baihan.huang@gmail.com>
Co-authored-by: Sherlock Huang <bahuang@OrtTrainingDev3.af05slrtruoetgaxwwjv5nsq5e.px.internal.cloudapp.net>
Co-authored-by: baijumeswani <bmeswani@microsoft.com>
Co-authored-by: Tixxx <tix@microsoft.com>
Co-authored-by: liqunfu <liqfu@microsoft.com>
Co-authored-by: Yulong Wang <yulongw@microsoft.com>
Co-authored-by: Guoyu Wang <62914304+gwang-msft@users.noreply.github.com>
Co-authored-by: Tianlei Wu <tlwu@microsoft.com>
This commit is contained in:
Xueyun Zhu 2021-05-28 09:10:40 -07:00 committed by GitHub
parent c487824a31
commit a54716e5a9
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
36 changed files with 1125 additions and 310 deletions

View file

@ -14,6 +14,25 @@ if (${CMAKE_SYSTEM_NAME} STREQUAL "iOS")
set(OUTPUT_STYLE xcode)
endif()
# This macro is to get the path of header files for mobile packaging, for iOS and Android
macro(get_mobile_api_headers _HEADERS)
# include both c and cxx api
set(${_HEADERS}
"${REPO_ROOT}/include/onnxruntime/core/session/onnxruntime_c_api.h"
"${REPO_ROOT}/include/onnxruntime/core/session/onnxruntime_cxx_api.h"
"${REPO_ROOT}/include/onnxruntime/core/session/onnxruntime_cxx_inline.h"
)
# need to add header files for enabled EPs
foreach(f ${ONNXRUNTIME_PROVIDER_NAMES})
file(GLOB _provider_headers CONFIGURE_DEPENDS
"${REPO_ROOT}/include/onnxruntime/core/providers/${f}/*.h"
)
list(APPEND ${_HEADERS} "${_provider_headers}")
unset(_provider_headers)
endforeach()
endmacro()
#If you want to verify if there is any extra line in symbols.txt, run
# nm -C -g --defined libonnxruntime.so |grep -v '\sA\s' | cut -f 3 -d ' ' | sort
# after build
@ -39,21 +58,7 @@ if(WIN32)
"${ONNXRUNTIME_ROOT}/core/dll/onnxruntime.rc"
)
elseif(onnxruntime_BUILD_APPLE_FRAMEWORK)
# include both c and cxx api
set(APPLE_FRAMEWORK_HEADERS
"${REPO_ROOT}/include/onnxruntime/core/session/onnxruntime_c_api.h"
"${REPO_ROOT}/include/onnxruntime/core/session/onnxruntime_cxx_api.h"
"${REPO_ROOT}/include/onnxruntime/core/session/onnxruntime_cxx_inline.h"
)
# need to add header files for enabled EPs
foreach(f ${ONNXRUNTIME_PROVIDER_NAMES})
file(GLOB _provider_headers CONFIGURE_DEPENDS
"${REPO_ROOT}/include/onnxruntime/core/providers/${f}/*.h"
)
list(APPEND APPLE_FRAMEWORK_HEADERS "${_provider_headers}")
unset(_provider_headers)
endforeach()
get_mobile_api_headers(APPLE_FRAMEWORK_HEADERS)
# apple framework requires the header file be part of the library
onnxruntime_add_shared_library(onnxruntime
@ -132,6 +137,18 @@ if(CMAKE_SYSTEM_NAME STREQUAL "Android" OR (onnxruntime_MINIMAL_BUILD AND UNIX))
endif()
endif()
# we need to copy C/C++ API headers to be packed into Android AAR package
if(CMAKE_SYSTEM_NAME STREQUAL "Android" AND onnxruntime_BUILD_JAVA)
get_mobile_api_headers(ANDROID_AAR_HEADERS)
set(ANDROID_HEADERS_DIR ${CMAKE_CURRENT_BINARY_DIR}/android/headers)
file(MAKE_DIRECTORY ${ANDROID_HEADERS_DIR})
# copy the header files one by one
foreach(h_ ${ANDROID_AAR_HEADERS})
get_filename_component(HEADER_NAME_ ${h_} NAME)
configure_file(${h_} ${ANDROID_HEADERS_DIR}/${HEADER_NAME_} COPYONLY)
endforeach()
endif()
target_link_libraries(onnxruntime PRIVATE
onnxruntime_session
${onnxruntime_libs}

View file

@ -7,7 +7,7 @@ This directory contains multiple NPM projects:
- [onnxruntime-web](#onnxruntime-web)
- [onnxruntime-react-native](#onnxruntime-react-native)
### Development
## Development
This folder contains a `.vscode` folder for Visual Studio Code workspace configs. Using VSCode to open this folder
will allow code-formatting and linting features on typescript and C/C++ source code inside this folder. Following files
@ -19,20 +19,45 @@ are used for code-formatting and linting features for developers:
- .eslintrc.js
- .clang-format
#### Using VSCode:
Please follow the steps described below to setup development environment.
1. in `<ORT_ROOT>/js`, run:
> npm ci
2. use VSCode to open folder `<ORT_ROOT>/js`
3. install VSCode extension if not installed yet:
- Clang-Format
- ESLint
### Prerequisites
To populate typescript type declarations, in each projects, run `npm ci`.
- Node.js (14.0+): https://nodejs.org/ - (Optional) Use nvm ([Windows](https://github.com/coreybutler/nvm-windows) / [Mac/Linux](https://github.com/creationix/nvm)) to install Node.js
#### Run code formatter and linter manually
- Python (2.7 or 3.6+): https://www.python.org/downloads/
in `<ORT_ROOT>/js`, use `npm run lint` to run ESLint , and use `npm run format` to run clang-format.
- python should be added to the PATH environment variable
- Visual Studio Code: https://code.visualstudio.com/
- **required** extension: [ESLint](https://marketplace.visualstudio.com/items?itemName=dbaeumer.vscode-eslint)
- **required** extension: [Clang-Format](https://marketplace.visualstudio.com/items?itemName=xaver.clang-format)
- **required** extension: [Debugger for Chrome](https://marketplace.visualstudio.com/items?itemName=msjsdiag.debugger-for-chrome)
- Chrome or Edge Browser
### Setup TypeScript development environment
In `<ORT_ROOT>/js`, run:
```
npm ci
```
This will install Clang-format and ESLint for code-formatting and linting features. This is a one-time setup unless a `git clean` is performed or folder `<ORT_ROOT>/js/node_modules` is removed manually.
### Using VSCode:
Use VSCode to open folder `<ORT_ROOT>/js`.
Make sure to open the correct folder to allow VSCode to load workspace configuration. Otherwise typescript and code formatter may not work as expected.
To populate typescript type declarations, in each project folder, run `npm ci`.
### Run code formatter and linter manually
In `<ORT_ROOT>/js`, use `npm run lint` to run ESLint , and use `npm run format` to run clang-format.
## onnxruntime-common
@ -112,10 +137,6 @@ It should be able to consumed by from projects that uses NPM packages (through a
This project is a library for running ONNX models on browsers. It is the successor of [ONNX.js](https://github.com/Microsoft/onnxjs).
### Requirements
Node.js v12+ (recommended v14+)
### Build
1. Install NPM packages
@ -124,9 +145,9 @@ Node.js v12+ (recommended v14+)
2. in `<ORT_ROOT>/js/common/`, run `npm ci`.
3. in `<ORT_ROOT>/js/web/`, run `npm ci`.
2. ~~Follow [instructions](https://www.onnxruntime.ai/docs/how-to/build.html#apis-and-language-bindings) for building ONNX Runtime WebAssembly. (TODO: document is not ready. we are working on it.)~~
2. ~~Follow [instructions](https://www.onnxruntime.ai/docs/how-to/build.html#apis-and-language-bindings) for building ONNX Runtime WebAssembly. (TODO: document is not ready. we are working on it. Please see steps described as below.)~~
in `<ORT_ROOT>/`, run either of the following commands to build WebAssembly:
in `<ORT_ROOT>/`, run one of the following commands to build WebAssembly:
```sh
# In windows, use 'build' to replace './build.sh'
@ -134,11 +155,16 @@ Node.js v12+ (recommended v14+)
# The following command build debug.
./build.sh --build_wasm
# The following command build debug with debug info.
./build.sh --build_wasm --skip_tests --enable_wasm_debug_info
# The following command build release.
./build.sh --config Release --build_wasm --skip_tests --disable_wasm_exception_catching --disable_rtti
```
To build with multi-thread support, append flag ` --enable_wasm_threads` to the command. Make sure to build both single-thread and multi-thread before next step.
To build with multi-thread support, append flag `--enable_wasm_threads` to the command. Make sure to build both single-thread and multi-thread before next step.
NOTE: You can also find latest build artifacts on [Windows WebAssembly CI Pipeline](https://dev.azure.com/onnxruntime/onnxruntime/_build?definitionId=161&_a=summary&repositoryFilter=1&branchFilter=4%2C4%2C4%2C4%2C4%2C4). Choose any build for master branch, download artifacts "Release_ort-wasm" and "Release_ort-wasm-threaded" and unzip.
3. Copy following files from build output folder to `<ORT_ROOT>/js/web/dist/`:
@ -156,9 +182,91 @@ Node.js v12+ (recommended v14+)
npm run build
```
### Test
We use command `npm test` (test runner) and `npm run test:e2e` (E2E test) for tests in ONNXRuntime Web.
#### test runner
In folder `<ORT_ROOT>/js/web`,
- Run `npm test -- --help` for a full CLI instruction.
- Run `npm test -- <your-args> --debug` to run one or more test cases.
There are multiple levels of tests for ONNXRuntime Web:
- unit test: tests for individual components written in TypeScript. Launch unit test by:
```
npm test -- unittest
```
- model test: run a single model. The model folder should contains one .onnx model file and one or more folders for test cases, each folder contains several input*\*.pb and output*\*.pb as test data. Launch model test by:
```
npm test -- model <model_folder>
```
- op test: test a single operator. An op test is described in a `.jsonc` file which specify the operator type, its attributes and one or more test case(s), each includes a list of expected input tensor(s) and output tensor(s). The `.jsonc` file is located at `<ORT_ROOT>/js/web/test/data/ops`. Launch op test by:
```
npm test -- op <file_name>
```
- suite test: suite test includes unit test, a list of model tests and op tests. Launch suite test by:
```
npm test
```
#### E2E test
E2E test is for testing end-to-end package consuming. In this test, NPM packages for `onnxruntime-common` and `onnxruntime-web` are generated and a clean folder is used for installing packages. Then a simple mocha test is performed to make sure package can be consumed correctly.
To launch E2E test:
```
npm run test:e2e
```
### Debugging
#### Debugging TypeScript on Desktop/Chrome
To debug the code from test-runner on Chrome:
- Launch `npm test -- <your_args> --debug`. It opens an instance of Chrome browser.
- In the open Chrome browser, click the `DEBUG` button on the top-right of the page.
- In VSCode, click [side bar]->Run and Debug->select [Attach to Chrome]->click [Start Debugging] to attach.
- put breakpoints in source code, and Refresh the page to reload.
#### Debugging TypeScript on iOS/Safari
To debug on an Apple iOS device, please refer to the following steps:
- install [
RemoteDebug iOS WebKit Adapter](https://github.com/RemoteDebug/remotedebug-ios-webkit-adapter) by following its instructions.
- launch the adapter in commandline: `remotedebug_ios_webkit_adapter --port=9000`.
- in VSCode, select debug configuration `Remote Browser via Webkit Adaptor`.
- follow the steps above to debug.
#### Debugging TypeScript on Android/Chrome
To debug on an Android device, please refer to the following steps:
- Install [Android SDK Platform Tools](https://developer.android.com/studio/releases/platform-tools) and make sure `adb` is ready to use.
- Follow instructions in [Remote Debugging on Android](https://developer.chrome.com/devtools/docs/remote-debugging-legacy) to launch `adb`. Make sure to use port 9000 so that the existing debug configuration works.
- in VSCode, select debug configuration `Remote Browser via Webkit Adaptor`.
- follow the steps above to debug.
#### Debugging C/C++ for ONNX Runtime WebAssembly
To debug C/C++ code for ONNX Runtime WebAssembly, you need to build ONNX Runtime with debug info (see [Build](#Build-2)).
Currently debugging C/C++ code in WebAssembly is not supported in VSCode yet. Please follow [this instruction](https://developer.chrome.com/blog/wasm-debugging-2020/) to debug in browser devtool using extension [C/C++ DevTools Support (DWARF)](https://chrome.google.com/webstore/detail/cc%20%20-devtools-support-dwa/pdcpmagijalfljmkmjngeonclgbbannb).
### Generating Document
Use command `npm run build:doc` to generate the latest documents.
### Distribution
It should be able to consumed by both from projects that uses NPM packages (through a Node.js folder structure of `node_modules` folder that generated by `npm install onnxruntime-web`) and from a CDN service that serves a `.min.js` file and one or multiple `.wasm` file(s).
It should be able to consumed by both from projects that uses NPM packages (through a Node.js folder structure of `node_modules` folder that generated by `npm install onnxruntime-web`) and from a CDN service that serves a `ort.min.js` file and one or multiple `.wasm` file(s).
## onnxruntime-react-native
@ -192,6 +300,7 @@ This project provides an ONNX Runtime React Native JavaScript library to run ONN
1. Set up an Android build environment referring to [instruction](https://www.onnxruntime.ai/docs/how-to/build.html#android)
2. In `<ORT_ROOT>`, run this python script to build ONNX Runtime Android archive file. In windows, this requires admin account to build. If an app uses a fixed set of models, refer to [instruction](https://www.onnxruntime.ai/docs/how-to/build.html#android) and build a mobile version package
```python
python tools/ci_build/github/android/build_aar_package.py js/react_native/scripts/aar_build_settings.json --config MinSizeRel --android_sdk_path <ANDROID_SDK_PATH> --android_ndk_path <ANDROID_NDK_PATH> --build_dir <BUILD_DIRECTORY>
```
@ -199,6 +308,7 @@ This project provides an ONNX Runtime React Native JavaScript library to run ONN
3. This generates `onnxruntime-mobile-<version>.aar` in `<BUILD_DIRECTORY>/aar_out/MinSizeRel/com/microsoft/onnxruntime/onnxruntime-mobile/<version>`. Copy `aar` file into `<ORT_ROOT>/js/react_native/android/libs` and rename it as `onnxruntime.aar`
4. To verify, open Android Emulator and run this command from `<ORT_ROOT>/js/react_native/android`
```sh
adb shell am instrument -w ai.onnxruntime.react_native.test/androidx.test.runner.AndroidJUnitRunner
```
@ -206,40 +316,45 @@ This project provides an ONNX Runtime React Native JavaScript library to run ONN
3. Build iOS ONNX Runtime package
1. Set up iOS build environment referring to [instruction](https://www.onnxruntime.ai/docs/how-to/build.html#ios).
2. Build ONNX Runtime library for iOS from `<ORT_ROOT>` using this command,
```sh
./build.sh --config MinSizeRel --use_xcode --ios --ios_sysroot iphoneos --osx_arch arm64 --apple_deploy_target 11
```
Copy `<ORT_ROOT>/build/iOS/MinSizeRel/MinSizeRel-iphoneos/libonnxruntime.<version>.dylib` file into `<ORT_ROOT>/js/react_native/ios/Libraries/onnxruntime/lib/iphoneos`
3. Clean up the previous build and build ONNX Runtime library for iOS Simulator from `<ORT_ROOT>`
```sh
./build.sh --config MinSizeRel --use_xcode --ios --ios_sysroot iphonesimulator --osx_arch x86_64 --apple_deploy_target 11
```
Copy `<ORT_ROOT>/build/iOS/MinSizeRel/MinSizeRel-iphonesimulator/libonnxruntime.<version>.dylib` file into `<ORT_ROOT>/js/react_native/ios/Libraries/onnxruntime/lib/iphonesimulator`
4. Edit `onnxruntime-react-native.iphoneos.podspec` and `onnxruntime-react-native.iphonesimulator.podsepc` in `<ORT_ROOT>/js/react_native` to change a version of ONNX Runtime library.
5. Copy ONNX Runtime header files
```sh
cp <ORT_ROOT>/include/onnxruntime/core/session/*.h <ORT_ROOT>/js/react_native/ios/Libraries/onnxruntime/include
```
6. To verify, open iOS Simulator and run this command from `<ORT_ROOT>/js/react_native/ios`. Change a destination to specify a running iOS Simulator.
```sh
pod install
export ONNXRUNTIME_VERSION=<version>; xcodebuild test -workspace OnnxruntimeModule.xcworkspace -scheme OnnxruntimeModuleTest -destination 'platform=iOS Simulator,name=iPhone 11,OS=14.5'
```
```sh
pod install
export ONNXRUNTIME_VERSION=<version>; xcodebuild test -workspace OnnxruntimeModule.xcworkspace -scheme OnnxruntimeModuleTest -destination 'platform=iOS Simulator,name=iPhone 11,OS=14.5'
```
4. Update a version in `package.json` to align with ONNX Runtime version.
5. Test an example for Android and iOS. In Windows, open Android Emulator first. From `<ORT_ROOT>/js/react_native`
```sh
yarn bootstrap
yarn example ios
yarn example android
```
```sh
yarn bootstrap
yarn example ios
yarn example android
```
### NPM Packaging
@ -251,7 +366,7 @@ This project provides an ONNX Runtime React Native JavaScript library to run ONN
4. Run `npm publish <tgz> --dry-run` to see how it's going to be published
5. Run `npm publish <tgz>` to publish to npmjs
5. Run `npm publish <tgz>` to publish to npmjs. If it's for a dev, add flag `--tag dev`.
### Distribution

View file

@ -6,7 +6,7 @@ ONNX Runtime Web has adopted WebAssembly and WebGL technologies for providing an
### Why ONNX models
The [Open Neural Network Exchange](http://onnx.ai/) (ONNX) is an open standard for representing machine learning models. The biggest advantage of ONNX is that it allows interoperability across different open source AI frameworks, which itself offers more flexibility for AI frameworks adoption. See [Getting ONNX Models](#Getting-ONNX-models).
The [Open Neural Network Exchange](http://onnx.ai/) (ONNX) is an open standard for representing machine learning models. The biggest advantage of ONNX is that it allows interoperability across different open source AI frameworks, which itself offers more flexibility for AI frameworks adoption.
### Why ONNX Runtime Web
@ -22,35 +22,25 @@ Refer to [ONNX Runtime JavaScript examples](https://github.com/microsoft/onnxrun
## Documents
### Developers
### Developement
Refer to [Using VSCode](../README.md#Using-VSCode) for setting up development environment.
Refer to the following links for development information:
For information about building ONNX Runtime Web development, please check [Build](../README.md#build-2).
### Getting ONNX models
You can get ONNX models easily in multiple ways:
- Choose a pre-trained ONNX model from the [ONNX Model Zoo](https://github.com/onnx/models)
- Convert models from mainstream frameworks, e.g. PyTorch, TensorFlow and Keras, by following [ONNX tutorials](https://github.com/onnx/tutorials)
- Use your data to generate a customized ONNX model from [Azure Custom Vision service](https://docs.microsoft.com/en-us/azure/cognitive-services/Custom-Vision-Service/home)
- [Train a custom model in AzureML](https://github.com/Azure/MachineLearningNotebooks/tree/master/training) and save it in the ONNX format
Learn more about ONNX
- [ONNX website](http://onnx.ai/)
- [ONNX on GitHub](https://github.com/onnx/onnx)
- [Development](../README.md#Development)
- [Build](../README.md#Build-2)
- [Test](../README.md#Test)
- [Debugging](../README.md#Debugging)
- [Generating Document](../README.md#Generating-Document)
### Compatibility
| OS/Browser | Chrome | Edge | Safari | Electron |
| :--------------: | :----------------: | :----------------: | :----------------: | :----------------: |
| Windows 10 | :heavy_check_mark: | :heavy_check_mark: | - | :heavy_check_mark: |
| macOS | :heavy_check_mark: | - | :heavy_check_mark: | :heavy_check_mark: |
| Ubuntu LTS 18.04 | :heavy_check_mark: | - | - | :heavy_check_mark: |
| iOS | :heavy_check_mark: | :heavy_check_mark: | :heavy_check_mark: | - |
| Android | :heavy_check_mark: | - | - | - |
| OS/Browser | Chrome | Edge | Safari | Electron | Node.js |
| :--------------: | :---------: | :---------: | :----: | :---------: | :-----: |
| Windows 10 | wasm, webgl | wasm, webgl | - | wasm, webgl | wasm |
| macOS | wasm | - | wasm | wasm | wasm |
| Ubuntu LTS 18.04 | wasm | - | - | wasm | wasm |
| iOS | wasm | wasm | wasm | - | - |
| Android | wasm | - | - | - | - |
### Operators
@ -60,7 +50,7 @@ ONNX Runtime Web currently support all operators in [ai.onnx](https://github.com
#### WebGL backend
ONNX Runtime Web currently supports most operators in [ai.onnx](https://github.com/onnx/onnx/blob/rel-1.2.3/docs/Operators.md) operator set v7 (opset v7). See [operators.md](./docs/operators.md) for a complete, detailed list of which ONNX operators are supported by WebGL backend.
ONNX Runtime Web currently supports a subset of operators in [ai.onnx](https://github.com/onnx/onnx/blob/master/docs/Operators.md) operator set. See [operators.md](./docs/operators.md) for a complete, detailed list of which ONNX operators are supported by WebGL backend.
## License

174
js/web/docs/operators.md Normal file
View file

@ -0,0 +1,174 @@
## Operators Support Table
The following table shows [ai.onnx](https://github.com/onnx/onnx/blob/master/docs/Operators.md) operators from which onnx opset version are currently supported by onnxjs. For example, `4-6, 8+` means ONNX Runtime Web currently support opset version 4 to 6, 8 and above.
See [Compatibility](../README.md#Compatibility) for a list of the supported platforms.
*This file is automatically generated from the def files via [this script](../script/generate-operator-md.ts). Do not modify directly.*
| Operator | WebGl Backend |
|:--------:|:-------------:|
| [Abs](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Abs) | [6-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Abs-6), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Abs-13) |
| [Acos](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Acos) | [7+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Acos-7) |
| [Acosh](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Acosh) | |
| [Add](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Add) | [7-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Add-7), [13](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Add-13), [14+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Add-14) |
| [And](https://github.com/onnx/onnx/blob/master/docs/Operators.md#And) | [7+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#And-7) |
| [ArgMax](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ArgMax) | |
| [ArgMin](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ArgMin) | |
| [Asin](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Asin) | [7+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Asin-7) |
| [Asinh](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Asinh) | |
| [Atan](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Atan) | [7+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Atan-7) |
| [Atanh](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Atanh) | |
| [AveragePool](https://github.com/onnx/onnx/blob/master/docs/Operators.md#AveragePool) | [7-9](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#AveragePool-7), [10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#AveragePool-10) |
| [BatchNormalization](https://github.com/onnx/onnx/blob/master/docs/Operators.md#BatchNormalization) | [7-8](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#BatchNormalization-7), [9-13](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#BatchNormalization-9), [14+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#BatchNormalization-14) |
| [BitShift](https://github.com/onnx/onnx/blob/master/docs/Operators.md#BitShift) | |
| [Cast](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Cast) | |
| [Ceil](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Ceil) | [6-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Ceil-6), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Ceil-13) |
| [Celu](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Celu) | |
| [Clip](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Clip) | [6-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Clip-6) |
| [Compress](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Compress) | |
| [Concat](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Concat) | [4-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Concat-4), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Concat-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Concat-13) |
| [ConcatFromSequence](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ConcatFromSequence) | |
| [Constant](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Constant) | |
| [ConstantOfShape](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ConstantOfShape) | |
| [Conv](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Conv) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Conv-1), [11+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Conv-11) |
| [ConvInteger](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ConvInteger) | |
| [ConvTranspose](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ConvTranspose) | |
| [Cos](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Cos) | [7+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Cos-7) |
| [Cosh](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Cosh) | |
| [CumSum](https://github.com/onnx/onnx/blob/master/docs/Operators.md#CumSum) | |
| [DepthToSpace](https://github.com/onnx/onnx/blob/master/docs/Operators.md#DepthToSpace) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#DepthToSpace-1), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#DepthToSpace-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#DepthToSpace-13) |
| [DequantizeLinear](https://github.com/onnx/onnx/blob/master/docs/Operators.md#DequantizeLinear) | |
| [Det](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Det) | |
| [Div](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Div) | [7-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Div-7), [13](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Div-13), [14+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Div-14) |
| [Dropout](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Dropout) | [7-9](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Dropout-7), [10-11](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Dropout-10), [12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Dropout-12), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Dropout-13) |
| [DynamicQuantizeLinear](https://github.com/onnx/onnx/blob/master/docs/Operators.md#DynamicQuantizeLinear) | |
| [Einsum](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Einsum) | |
| [Elu](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Elu) | [6+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Elu-6) |
| [Equal](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Equal) | [7-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Equal-7), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Equal-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Equal-13) |
| [Erf](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Erf) | |
| [Exp](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Exp) | [6-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Exp-6), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Exp-13) |
| [Expand](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Expand) | |
| [EyeLike](https://github.com/onnx/onnx/blob/master/docs/Operators.md#EyeLike) | |
| [Flatten](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Flatten) | [1-8](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Flatten-1), [9-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Flatten-9), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Flatten-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Flatten-13) |
| [Floor](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Floor) | [6-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Floor-6), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Floor-13) |
| [GRU](https://github.com/onnx/onnx/blob/master/docs/Operators.md#GRU) | |
| [Gather](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Gather) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Gather-1), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Gather-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Gather-13) |
| [GatherElements](https://github.com/onnx/onnx/blob/master/docs/Operators.md#GatherElements) | |
| [GatherND](https://github.com/onnx/onnx/blob/master/docs/Operators.md#GatherND) | |
| [Gemm](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Gemm) | [7-8](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Gemm-7), [9-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Gemm-9), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Gemm-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Gemm-13) |
| [GlobalAveragePool](https://github.com/onnx/onnx/blob/master/docs/Operators.md#GlobalAveragePool) | [1+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#GlobalAveragePool-1) |
| [GlobalLpPool](https://github.com/onnx/onnx/blob/master/docs/Operators.md#GlobalLpPool) | |
| [GlobalMaxPool](https://github.com/onnx/onnx/blob/master/docs/Operators.md#GlobalMaxPool) | [1+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#GlobalMaxPool-1) |
| [Greater](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Greater) | [7-8](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Greater-7), [9-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Greater-9), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Greater-13) |
| [GreaterOrEqual](https://github.com/onnx/onnx/blob/master/docs/Operators.md#GreaterOrEqual) | |
| [HardSigmoid](https://github.com/onnx/onnx/blob/master/docs/Operators.md#HardSigmoid) | |
| [HardSwish](https://github.com/onnx/onnx/blob/master/docs/Operators.md#HardSwish) | |
| [Hardmax](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Hardmax) | |
| [Identity](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Identity) | [1-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Identity-1), [13](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Identity-13), [14+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Identity-14) |
| [If](https://github.com/onnx/onnx/blob/master/docs/Operators.md#If) | |
| [InstanceNormalization](https://github.com/onnx/onnx/blob/master/docs/Operators.md#InstanceNormalization) | [6+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#InstanceNormalization-6) |
| [IsInf](https://github.com/onnx/onnx/blob/master/docs/Operators.md#IsInf) | |
| [IsNaN](https://github.com/onnx/onnx/blob/master/docs/Operators.md#IsNaN) | |
| [LRN](https://github.com/onnx/onnx/blob/master/docs/Operators.md#LRN) | |
| [LSTM](https://github.com/onnx/onnx/blob/master/docs/Operators.md#LSTM) | |
| [LeakyRelu](https://github.com/onnx/onnx/blob/master/docs/Operators.md#LeakyRelu) | [6+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#LeakyRelu-6) |
| [Less](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Less) | [7-8](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Less-7), [9-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Less-9), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Less-13) |
| [LessOrEqual](https://github.com/onnx/onnx/blob/master/docs/Operators.md#LessOrEqual) | |
| [Log](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Log) | [6-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Log-6), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Log-13) |
| [LogSoftmax](https://github.com/onnx/onnx/blob/master/docs/Operators.md#LogSoftmax) | |
| [Loop](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Loop) | |
| [LpNormalization](https://github.com/onnx/onnx/blob/master/docs/Operators.md#LpNormalization) | |
| [LpPool](https://github.com/onnx/onnx/blob/master/docs/Operators.md#LpPool) | |
| [MatMul](https://github.com/onnx/onnx/blob/master/docs/Operators.md#MatMul) | [1-8](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#MatMul-1), [9-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#MatMul-9), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#MatMul-13) |
| [MatMulInteger](https://github.com/onnx/onnx/blob/master/docs/Operators.md#MatMulInteger) | |
| [Max](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Max) | |
| [MaxPool](https://github.com/onnx/onnx/blob/master/docs/Operators.md#MaxPool) | [1-7](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#MaxPool-1), [8-9](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#MaxPool-8) |
| [MaxRoiPool](https://github.com/onnx/onnx/blob/master/docs/Operators.md#MaxRoiPool) | |
| [MaxUnpool](https://github.com/onnx/onnx/blob/master/docs/Operators.md#MaxUnpool) | |
| [Mean](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Mean) | |
| [MeanVarianceNormalization](https://github.com/onnx/onnx/blob/master/docs/Operators.md#MeanVarianceNormalization) | |
| [Min](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Min) | |
| [Mod](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Mod) | |
| [Mul](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Mul) | [7-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Mul-7), [13](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Mul-13), [14+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Mul-14) |
| [Multinomial](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Multinomial) | |
| [Neg](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Neg) | [6-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Neg-6), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Neg-13) |
| [NegativeLogLikelihoodLoss](https://github.com/onnx/onnx/blob/master/docs/Operators.md#NegativeLogLikelihoodLoss) | |
| [NonMaxSuppression](https://github.com/onnx/onnx/blob/master/docs/Operators.md#NonMaxSuppression) | |
| [NonZero](https://github.com/onnx/onnx/blob/master/docs/Operators.md#NonZero) | |
| [Not](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Not) | [1+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Not-1) |
| [OneHot](https://github.com/onnx/onnx/blob/master/docs/Operators.md#OneHot) | |
| [Or](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Or) | [7+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Or-7) |
| [PRelu](https://github.com/onnx/onnx/blob/master/docs/Operators.md#PRelu) | [7-8](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#PRelu-7), [9+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#PRelu-9) |
| [Pad](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Pad) | [2-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Pad-2) |
| [Pow](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Pow) | [7-11](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Pow-7), [12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Pow-12), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Pow-13) |
| [QLinearConv](https://github.com/onnx/onnx/blob/master/docs/Operators.md#QLinearConv) | |
| [QLinearMatMul](https://github.com/onnx/onnx/blob/master/docs/Operators.md#QLinearMatMul) | |
| [QuantizeLinear](https://github.com/onnx/onnx/blob/master/docs/Operators.md#QuantizeLinear) | |
| [RNN](https://github.com/onnx/onnx/blob/master/docs/Operators.md#RNN) | |
| [RandomNormal](https://github.com/onnx/onnx/blob/master/docs/Operators.md#RandomNormal) | |
| [RandomNormalLike](https://github.com/onnx/onnx/blob/master/docs/Operators.md#RandomNormalLike) | |
| [RandomUniform](https://github.com/onnx/onnx/blob/master/docs/Operators.md#RandomUniform) | |
| [RandomUniformLike](https://github.com/onnx/onnx/blob/master/docs/Operators.md#RandomUniformLike) | |
| [Range](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Range) | |
| [Reciprocal](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Reciprocal) | |
| [ReduceL1](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ReduceL1) | |
| [ReduceL2](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ReduceL2) | |
| [ReduceLogSum](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ReduceLogSum) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceLogSum-1), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceLogSum-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceLogSum-13) |
| [ReduceLogSumExp](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ReduceLogSumExp) | |
| [ReduceMax](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ReduceMax) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceMax-1), [11](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceMax-11), [12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceMax-12), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceMax-13) |
| [ReduceMean](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ReduceMean) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceMean-1), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceMean-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceMean-13) |
| [ReduceMin](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ReduceMin) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceMin-1), [11](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceMin-11), [12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceMin-12), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceMin-13) |
| [ReduceProd](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ReduceProd) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceProd-1), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceProd-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceProd-13) |
| [ReduceSum](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ReduceSum) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceSum-1), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceSum-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceSum-13) |
| [ReduceSumSquare](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ReduceSumSquare) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceSumSquare-1), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceSumSquare-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#ReduceSumSquare-13) |
| [Relu](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Relu) | [6-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Relu-6), [13](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Relu-13), [14+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Relu-14) |
| [Reshape](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Reshape) | [5-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Reshape-5), [13](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Reshape-13), [14+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Reshape-14) |
| [Resize](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Resize) | [10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Resize-10), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Resize-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Resize-13) |
| [ReverseSequence](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ReverseSequence) | |
| [RoiAlign](https://github.com/onnx/onnx/blob/master/docs/Operators.md#RoiAlign) | |
| [Round](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Round) | |
| [Scan](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Scan) | |
| [Scatter](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Scatter) | |
| [ScatterElements](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ScatterElements) | |
| [ScatterND](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ScatterND) | |
| [Selu](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Selu) | |
| [SequenceAt](https://github.com/onnx/onnx/blob/master/docs/Operators.md#SequenceAt) | |
| [SequenceConstruct](https://github.com/onnx/onnx/blob/master/docs/Operators.md#SequenceConstruct) | |
| [SequenceEmpty](https://github.com/onnx/onnx/blob/master/docs/Operators.md#SequenceEmpty) | |
| [SequenceErase](https://github.com/onnx/onnx/blob/master/docs/Operators.md#SequenceErase) | |
| [SequenceInsert](https://github.com/onnx/onnx/blob/master/docs/Operators.md#SequenceInsert) | |
| [SequenceLength](https://github.com/onnx/onnx/blob/master/docs/Operators.md#SequenceLength) | |
| [Shape](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Shape) | |
| [Shrink](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Shrink) | |
| [Sigmoid](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Sigmoid) | [6-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Sigmoid-6), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Sigmoid-13) |
| [Sign](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Sign) | |
| [Sin](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Sin) | [7+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Sin-7) |
| [Sinh](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Sinh) | |
| [Size](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Size) | |
| [Slice](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Slice) | [1-9](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Slice-1), [10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Slice-10), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Slice-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Slice-13) |
| [Softmax](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Softmax) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Softmax-1), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Softmax-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Softmax-13) |
| [SoftmaxCrossEntropyLoss](https://github.com/onnx/onnx/blob/master/docs/Operators.md#SoftmaxCrossEntropyLoss) | |
| [Softplus](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Softplus) | |
| [Softsign](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Softsign) | |
| [SpaceToDepth](https://github.com/onnx/onnx/blob/master/docs/Operators.md#SpaceToDepth) | |
| [Split](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Split) | [2-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Split-2), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Split-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Split-13) |
| [SplitToSequence](https://github.com/onnx/onnx/blob/master/docs/Operators.md#SplitToSequence) | |
| [Sqrt](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Sqrt) | [6-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Sqrt-6), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Sqrt-13) |
| [Squeeze](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Squeeze) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Squeeze-1), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Squeeze-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Squeeze-13) |
| [StringNormalizer](https://github.com/onnx/onnx/blob/master/docs/Operators.md#StringNormalizer) | |
| [Sub](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Sub) | [7-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Sub-7), [13](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Sub-13), [14+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Sub-14) |
| [Sum](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Sum) | [6-7](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Sum-6), [8-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Sum-8), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Sum-13) |
| [Tan](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Tan) | [7+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Tan-7) |
| [Tanh](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Tanh) | [6-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Tanh-6), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Tanh-13) |
| [TfIdfVectorizer](https://github.com/onnx/onnx/blob/master/docs/Operators.md#TfIdfVectorizer) | |
| [ThresholdedRelu](https://github.com/onnx/onnx/blob/master/docs/Operators.md#ThresholdedRelu) | |
| [Tile](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Tile) | [6-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Tile-6), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Tile-13) |
| [TopK](https://github.com/onnx/onnx/blob/master/docs/Operators.md#TopK) | |
| [Transpose](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Transpose) | [1-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Transpose-1), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Transpose-13) |
| [Trilu](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Trilu) | |
| [Unique](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Unique) | |
| [Unsqueeze](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Unsqueeze) | [1-10](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Unsqueeze-1), [11-12](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Unsqueeze-11), [13+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Unsqueeze-13) |
| [Upsample](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Upsample) | [7-8](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Upsample-7), [9](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Upsample-9) |
| [Where](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Where) | |
| [Xor](https://github.com/onnx/onnx/blob/master/docs/Operators.md#Xor) | [7+](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#Xor-7) |

View file

@ -654,7 +654,7 @@ export class CoordsGlslLib extends GlslLib {
if (inRank === 1 && !isInputScalar && !isOutputScalar) {
output = `
return vec4(outputValue.xy, outputValue.xy);
return vec4(outputValue.xx, outputValue.yy);
`;
} else if (isInputScalar && !isOutputScalar) {
if (outRank === 1) {
@ -1168,7 +1168,7 @@ export class CoordsGlslLib extends GlslLib {
return ${funcName}(${getSqueezedParams(params, keptDims)});
}
`;
return new GlslLibRoutine(source, ['coordinates.sampleTexture']);
return new GlslLibRoutine(source, ['coordinates.sampleTexture', 'coordinates.uvFromFlat']);
}
const texNumR = inputLayout.width;

View file

@ -10,6 +10,8 @@ import {WebGLInferenceHandler} from '../inference-handler';
import {ProgramInfo, RunData, WebGLOperator} from '../types';
export class WebGLBinaryOp extends BinaryOp implements WebGLOperator {
private usePackedTexture?: boolean;
constructor(
typeConstraint: readonly Tensor.DataType[], protected glslFunc: GlslValueFunction, opType?: string,
resultType?: Tensor.DataType) {
@ -19,14 +21,20 @@ export class WebGLBinaryOp extends BinaryOp implements WebGLOperator {
return inferenceHandler.run(this, inputs);
}
createProgramInfo(handler: WebGLInferenceHandler, inputs: Tensor[]): ProgramInfo {
const inputLayouts = handler.session.pack ?
const isBroadcast = !ShapeUtil.areEqual(inputs[0].dims, inputs[1].dims);
// TODO fix bcast in packed mode.
if (this.usePackedTexture === undefined) {
this.usePackedTexture = !isBroadcast && handler.session.pack;
}
const inputLayouts = this.usePackedTexture ?
inputs.map(t => handler.getOrCreateTextureLayout(t, 4, true, t.dims, true)) :
inputs.map(t => handler.getOrCreateTextureLayout(t));
const ouputLayout = handler.session.pack ?
const ouputLayout = this.usePackedTexture ?
handler.createTextureLayoutFromShape(inputs[0].dims, 4, inputs[0].dims, {isPacked: true, reverseWH: true}) :
handler.createTextureLayoutFromShape(inputs[0].dims);
const isBroadcast = !ShapeUtil.areEqual(inputs[0].dims, inputs[1].dims);
if (isBroadcast) {
const outputShape = BroadcastUtil.calcShape(inputs[0].dims, inputs[1].dims, false);
if (!outputShape) {
@ -48,7 +56,7 @@ export class WebGLBinaryOp extends BinaryOp implements WebGLOperator {
${bBcast}
return ${this.glslFunc.name}(_A(aindices), _B(bindices));
}`;
const outputLayout = handler.session.pack ?
const outputLayout = this.usePackedTexture ?
handler.createTextureLayoutFromShape(outputShape, 4, outputShape, {isPacked: true, reverseWH: true}) :
handler.createTextureLayoutFromShape(outputShape);
@ -57,8 +65,8 @@ export class WebGLBinaryOp extends BinaryOp implements WebGLOperator {
outputLayout,
samplers: ['A', 'B'],
shaderSource,
expectPackedInputs: handler.session.pack,
expectPackedOutputs: handler.session.pack
expectPackedInputs: this.usePackedTexture,
expectPackedOutputs: this.usePackedTexture
};
}
const glsl = getGlsl(handler.session.backend.glContext.version);
@ -71,7 +79,8 @@ export class WebGLBinaryOp extends BinaryOp implements WebGLOperator {
${glsl.output} = result;
}
`;
if (handler.session.pack) {
if (this.usePackedTexture) {
return {
hasMain: true,
inputLayouts,
@ -92,7 +101,7 @@ export class WebGLBinaryOp extends BinaryOp implements WebGLOperator {
}
}
createRunData(handler: WebGLInferenceHandler, programInfo: ProgramInfo, inputs: Tensor[]): RunData {
const inputTDs = handler.session.pack ?
const inputTDs = this.usePackedTexture ?
inputs.map((t) => handler.getOrCreateTextureData(t, handler.getOrCreateTextureLayout(t, 1, false, [], true))) :
inputs.map((t, i) => handler.getOrCreateTextureData(t, programInfo.inputLayouts[i]));
return {
@ -159,7 +168,7 @@ export function glslEqual(): GlslValueFunction {
return float(a == b);
}
vec4 ${name}(vec4 v1, vec4 v2) {
return vec4( v1 == v2 );
return vec4(equal(v1, v2));
}
`;
return {body, name, type: FunctionType.ValueBased};

View file

@ -47,11 +47,11 @@ export class WebGLIm2ColPacked implements WebGLOperator {
pos = rc.y + ${row};
if(blockIndex < ${im2colShape[1]} && pos < ${im2colShape[0]}) {
offsetY = int(blockIndex / (${this.convOutputShape[rank - 1]})) * ${this.strides[0]} - ${this.pads[1]};
offsetY = int(blockIndex / (${this.convOutputShape[rank - 1]})) * ${this.strides[0]} - ${this.pads[0]};
d0 = offsetY + ${this.dilations[0]} * (imod(pos, ${kernelSize}) / ${wshape[2]});
if(d0 < ${xshape[rowDim]} && d0 >= 0) {
offsetX = imod(blockIndex, ${this.convOutputShape[rank - 1]}) * ${this.strides[1]} - ${this.pads[0]};
offsetX = imod(blockIndex, ${this.convOutputShape[rank - 1]}) * ${this.strides[1]} - ${this.pads[1]};
d1 = offsetX + ${this.dilations[1]} * imod(imod(pos, ${kernelSize}), ${wshape[2]});
if(d1 < ${xshape[colDim]} && d1 >= 0) {

View file

@ -3,12 +3,14 @@
import {MatMul} from '../../../ops/matmul';
import {Tensor} from '../../../tensor';
import {BroadcastUtil} from '../../../util';
import {BroadcastUtil, ShapeUtil} from '../../../util';
import {WebGLInferenceHandler} from '../inference-handler';
import {ProgramInfo, RunData, WebGLOperator} from '../types';
import {WebGLMatMulPacked} from './matmul-pack';
export class WebGLMatMul extends MatMul implements WebGLOperator {
private usePackedTexture?: boolean;
packedImpl: WebGLMatMulPacked;
unpackedImpl: WebGLUnpackedMatMul;
constructor() {
@ -18,7 +20,12 @@ export class WebGLMatMul extends MatMul implements WebGLOperator {
}
run(inferenceHandler: WebGLInferenceHandler, inputs: Tensor[]): Tensor[] {
if (inferenceHandler.session.pack) {
if (this.usePackedTexture === undefined) {
const isBroadcast = !ShapeUtil.areEqual(inputs[0].dims, inputs[1].dims);
this.usePackedTexture = !isBroadcast && inferenceHandler.session.pack;
}
if (this.usePackedTexture) {
return inferenceHandler.run(this.packedImpl, inputs);
} else {
return inferenceHandler.run(this.unpackedImpl, inputs);
@ -26,7 +33,12 @@ export class WebGLMatMul extends MatMul implements WebGLOperator {
}
createProgramInfo(handler: WebGLInferenceHandler, inputs: Tensor[]): ProgramInfo {
if (handler.session.pack && inputs[0].dims.length > 1) {
if (this.usePackedTexture === undefined) {
const isBroadcast = !ShapeUtil.areEqual(inputs[0].dims, inputs[1].dims);
this.usePackedTexture = !isBroadcast && handler.session.pack;
}
if (this.usePackedTexture && inputs[0].dims.length > 1) {
return this.packedImpl.createProgramInfo(handler, inputs);
} else {
return this.unpackedImpl.createProgramInfo(handler, inputs);
@ -34,7 +46,7 @@ export class WebGLMatMul extends MatMul implements WebGLOperator {
}
createRunData(handler: WebGLInferenceHandler, programInfo: ProgramInfo, inputs: Tensor[]): RunData {
if (handler.session.pack && inputs[0].dims.length > 1) {
if (this.usePackedTexture && inputs[0].dims.length > 1) {
return this.packedImpl.createRunData(handler, programInfo, inputs);
} else {
return this.unpackedImpl.createRunData(handler, programInfo, inputs);

View file

@ -16,6 +16,7 @@
"scripts": {
"prepare": "tsc",
"build": "node ./script/build",
"build:doc": "node ./script/generate-operator-md",
"test": "node ./script/prepare-test-data && node ./script/test-runner-cli",
"test:e2e": "node ./test/e2e/run",
"prepack": "node ./script/prepack"

View file

@ -0,0 +1,105 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
import * as assert from 'assert';
import * as fs from 'fs';
import {EOL} from 'os';
import * as path from 'path';
import {Attribute} from '../lib/onnxjs/attribute';
import {WEBGL_OP_RESOLVE_RULES} from '../lib/onnxjs/backends/webgl/op-resolve-rules';
import {Operator} from '../lib/onnxjs/operators';
import {OpSet, resolveOperator} from '../lib/onnxjs/opset';
function checkSupport(type: string, range: [number, number], rules: readonly OpSet.ResolveRule[]) {
const node = {name: '', opType: type, inputs: [], outputs: [], attributes: new Attribute(undefined)};
for (let i = range[0]; i <= range[1]; i++) {
try {
resolveOperator(node, [{domain: '', version: i}], rules);
} catch (_e) {
return false;
}
}
return true;
}
function formatDesc(opType: string, range: [number, number], support: boolean, last: boolean) {
let versionDesc = '';
if (support) {
versionDesc = last ? `${range[0]}+` : range[0] === range[1] ? `${range[0]}` : `${range[0]}-${range[1]}`;
versionDesc = `[${versionDesc}](https://github.com/onnx/onnx/blob/master/docs/Changelog.md#${opType}-${range[0]})`;
}
return versionDesc;
}
function dummyOpConstructor(): Operator {
return {} as any as Operator;
}
const ops = new Map<string, Map<string, number[]>>();
const webglCheckOnlyRules =
WEBGL_OP_RESOLVE_RULES.map(rule => [rule[0], rule[1], rule[2], dummyOpConstructor] as OpSet.ResolveRule);
fs.readFileSync(path.join(__dirname, '../../../cmake/external/onnx/onnx/defs/operator_sets.h'), 'utf8')
.split(/\r?\n/)
.forEach(line => {
const matcher = /class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME\(\s*(\w+),\s*(\d+),\s*(\w+)\)/;
const matches = matcher.exec(line);
if (matches) {
const opset = matches[1];
const version = Number.parseInt(matches[2], 10);
const opType = matches[3];
let currentSet = ops.get(opset);
if (currentSet === undefined) {
currentSet = new Map<string, number[]>();
ops.set(opset, currentSet);
}
let currentOp = currentSet.get(opType);
if (currentOp === undefined) {
currentOp = [];
currentSet.set(opType, currentOp);
}
currentOp.push(version);
}
});
const opsets = Array.from(ops.keys());
assert.ok(opsets.length === 1 && opsets[0] === 'Onnx');
const onnxOpset = ops.get(opsets[0])!;
const opTypes = Array.from(onnxOpset.keys()).sort();
const doc = fs.createWriteStream(path.join(__dirname, '../docs/operators.md'));
doc.write(`## Operators Support Table${EOL}${EOL}`);
doc.write(`The following table shows [ai.onnx](https://github.com/onnx/onnx/blob/master/docs/Operators.md)\
operators from which onnx opset version are currently supported by onnxjs. For example, \`4-6, 8+\` means\
ONNX Runtime Web currently support opset version 4 to 6, 8 and above.${EOL}${EOL}`);
doc.write(`See [Compatibility](../README.md#Compatibility) for a list of the supported platforms.${EOL}${EOL}`);
doc.write(`*This file is automatically generated from the\
def files via [this script](../script/generate-operator-md.ts).\
Do not modify directly.*${EOL}${EOL}`);
doc.write(`| Operator | WebGl Backend |${EOL}`);
doc.write(`|:--------:|:-------------:|${EOL}`);
let VERSION_MAX = 0;
onnxOpset.forEach(versions => {
versions.forEach(version => VERSION_MAX = Math.max(VERSION_MAX, version));
});
for (const type of opTypes) {
const versions = onnxOpset.get(type)!.sort((a, b) => a - b);
const webgl: string[] = [];
for (let i = 0; i < versions.length; i++) {
const last = i === versions.length - 1;
const versionRange: [number, number] = [versions[i], last ? VERSION_MAX : versions[i + 1] - 1];
webgl.push(formatDesc(type, versionRange, checkSupport(type, versionRange, webglCheckOnlyRules), last));
}
doc.write(`| [${type}](https://github.com/onnx/onnx/blob/master/docs/Operators.md#${type}) | ${
webgl.filter(d => d.length > 0).join(', ')} |${EOL}`);
}
doc.end();

View file

@ -27,12 +27,12 @@
"test_averagepool_1d_default",
"test_averagepool_2d_default",
//"v12/test_averagepool_2d_pads", // TODO: fix avgpool and maxpool on VM
"v12/test_averagepool_2d_precomputed_pads",
"v12/test_averagepool_2d_precomputed_same_upper",
"v12/test_averagepool_2d_precomputed_strides",
"v12/test_averagepool_2d_same_upper",
"v12/test_averagepool_2d_same_lower",
"v12/test_averagepool_2d_strides",
// "v12/test_averagepool_2d_precomputed_pads",
// "v12/test_averagepool_2d_precomputed_same_upper",
// "v12/test_averagepool_2d_precomputed_strides",
// "v12/test_averagepool_2d_same_upper",
// "v12/test_averagepool_2d_same_lower",
// "v12/test_averagepool_2d_strides",
"test_averagepool_3d_default",
"test_basic_conv_with_padding",
"test_basic_conv_without_padding",
@ -102,13 +102,13 @@
"test_matmul_4d",
"test_maxpool_1d_default",
"test_maxpool_2d_default",
"v12/test_maxpool_2d_pads",
"v12/test_maxpool_2d_precomputed_pads",
"v12/test_maxpool_2d_precomputed_same_upper",
"v12/test_maxpool_2d_precomputed_strides",
"v12/test_maxpool_2d_same_lower",
"v12/test_maxpool_2d_same_upper",
"v12/test_maxpool_2d_strides",
// "v12/test_maxpool_2d_pads",
// "v12/test_maxpool_2d_precomputed_pads",
// "v12/test_maxpool_2d_precomputed_same_upper",
// "v12/test_maxpool_2d_precomputed_strides",
// "v12/test_maxpool_2d_same_lower",
// "v12/test_maxpool_2d_same_upper",
// "v12/test_maxpool_2d_strides",
"test_maxpool_3d_default",
"test_mul_bcast",
"test_mul_example",

View file

@ -2,10 +2,9 @@
// Licensed under the MIT License.
import {expect} from 'chai';
import {env} from 'onnxruntime-common';
import {Attribute} from '../../../../lib/onnxjs/attribute';
import {Backend, InferenceHandler, resolveBackend, SessionHandler} from '../../../../lib/onnxjs/backend';
import {WebGLBackend} from '../../../../lib/onnxjs/backends/backend-webgl';
import {WebGLInferenceHandler} from '../../../../lib/onnxjs/backends/webgl/inference-handler';
import {WebGLConcat} from '../../../../lib/onnxjs/backends/webgl/ops/concat';
import {Profiler} from '../../../../lib/onnxjs/instrument';
@ -207,17 +206,10 @@ describe('#UnitTest# - packed concat - Tensor concat', () => {
before('Initialize Context', async () => {
const profiler = Profiler.create();
backend = await resolveBackend('webgl');
// Explicitly set to true to trigger packed version
(backend as WebGLBackend).pack = true;
sessionhandler = backend.createSessionHandler({profiler});
inferenceHandler = sessionhandler.createInferenceHandler();
});
// Set it back to false, apparently this state is sticky throughout all the tests running in same browser session..
after('Resetting Context', () => {
(backend as WebGLBackend).pack = false;
});
const testDataSet = getTestData();
for (let k = 0; k < testDataSet.length; ++k) {
const testData = testDataSet[k];
@ -231,6 +223,11 @@ describe('#UnitTest# - packed concat - Tensor concat', () => {
return;
}
if (!env.webgl.pack) {
console.log('Skipping in unpacked texture mode.');
return;
}
const op = new WebGLConcat();
const attributes = new Attribute(undefined);
const axis = testData.axis;

View file

@ -2,10 +2,8 @@
// Licensed under the MIT License.
import {expect} from 'chai';
import {Attribute} from '../../../../lib/onnxjs/attribute';
import {Backend, InferenceHandler, resolveBackend, SessionHandler} from '../../../../lib/onnxjs/backend';
import {WebGLBackend} from '../../../../lib/onnxjs/backends/backend-webgl';
import {WebGLInferenceHandler} from '../../../../lib/onnxjs/backends/webgl/inference-handler';
import {WebGLDepthToSpace} from '../../../../lib/onnxjs/backends/webgl/ops/depth-to-space';
import {Profiler} from '../../../../lib/onnxjs/instrument';
@ -126,11 +124,6 @@ describe('#UnitTest# - unpacked WebGLDepthToSpace - Tensor WebGLDepthToSpace', (
inferenceHandler = sessionhandler.createInferenceHandler();
});
// Set it back to false, apparently this state is sticky throughout all the tests running in same browser session..
after('Resetting Context', () => {
(backend as WebGLBackend).pack = false;
});
const testDataSet = getTestData();
for (let k = 0; k < testDataSet.length; ++k) {
const testData = testDataSet[k];

View file

@ -2,9 +2,9 @@
// Licensed under the MIT License.
import {expect} from 'chai';
import {env} from 'onnxruntime-common';
import {Backend, InferenceHandler, resolveBackend, SessionHandler} from '../../../../lib/onnxjs/backend';
import {WebGLBackend} from '../../../../lib/onnxjs/backends/backend-webgl';
import {WebGLInferenceHandler} from '../../../../lib/onnxjs/backends/webgl/inference-handler';
import {WebGLMatMulPacked} from '../../../../lib/onnxjs/backends/webgl/ops/matmul-pack';
import {Profiler} from '../../../../lib/onnxjs/instrument';
@ -140,17 +140,10 @@ describe('#UnitTest# - packed matmul - Tensor matmul', () => {
before('Initialize Context', async () => {
const profiler = Profiler.create();
backend = await resolveBackend('webgl');
// Explicitly set to true to trigger packed version
(backend as WebGLBackend).pack = true;
sessionhandler = backend.createSessionHandler({profiler});
inferenceHandler = sessionhandler.createInferenceHandler();
});
// Set it back to false, apparently this state is sticky throughout all the tests running in same browser session..
after('Resetting Context', () => {
(backend as WebGLBackend).pack = false;
});
const testDataSet = getTestData();
for (let k = 0; k < testDataSet.length; ++k) {
const testData = testDataSet[k];
@ -164,6 +157,11 @@ describe('#UnitTest# - packed matmul - Tensor matmul', () => {
return;
}
if (!env.webgl.pack) {
console.log('Skipping in unpacked texture mode.');
return;
}
const op = new WebGLMatMulPacked();
const elementCountA = testData.elementCountA;

View file

@ -2,9 +2,8 @@
// Licensed under the MIT License.
import {expect} from 'chai';
import {env} from 'onnxruntime-common';
import {Backend, InferenceHandler, resolveBackend, SessionHandler} from '../../../../lib/onnxjs/backend';
import {WebGLBackend} from '../../../../lib/onnxjs/backends/backend-webgl';
import {WebGLInferenceHandler} from '../../../../lib/onnxjs/backends/webgl/inference-handler';
import {WebGLReshapePacked} from '../../../../lib/onnxjs/backends/webgl/ops/reshape-packed';
import {Profiler} from '../../../../lib/onnxjs/instrument';
@ -111,17 +110,10 @@ describe('#UnitTest# - reshape - packed', () => {
before('Initialize Context', async () => {
const profiler = Profiler.create();
backend = await resolveBackend('webgl');
// Explicitly set to true to trigger packed version
(backend as WebGLBackend).pack = true;
sessionhandler = backend.createSessionHandler({profiler});
inferenceHandler = sessionhandler.createInferenceHandler();
});
// Set it back to false, apparently this state is sticky throughout all the tests running in same browser session..
after('Resetting Context', () => {
(backend as WebGLBackend).pack = false;
});
const testDataSet = getTestData();
for (let k = 0; k < testDataSet.length; ++k) {
const testData = testDataSet[k];
@ -135,6 +127,11 @@ describe('#UnitTest# - reshape - packed', () => {
return;
}
if (!env.webgl.pack) {
console.log('Skipping in unpacked texture mode.');
return;
}
const op = new WebGLReshapePacked();
const elementCount = testData.elementCount;

View file

@ -54,8 +54,22 @@ ORT_SPECIFY_OP_KERNEL_ARG_DEFAULT_TYPES(
int8_t,
uint8_t);
ORT_SPECIFY_OP_KERNEL_ARG_DEFAULT_TYPES(
kCpuExecutionProvider, kOnnxDomain, Pad, 13, Input, 0,
float,
double,
int32_t,
int64_t,
uint32_t,
uint64_t,
int8_t,
uint8_t,
bool);
ORT_SPECIFY_OP_KERNEL_ARG_REQUIRED_TYPES(
kCpuExecutionProvider, kOnnxDomain, Pad, 11, Input, 0, int32_t, int64_t);
ORT_SPECIFY_OP_KERNEL_ARG_REQUIRED_TYPES(
kCpuExecutionProvider, kOnnxDomain, Pad, 13, Input, 0, int32_t, int64_t);
} // namespace op_kernel_type_control
using Pad2Types = ORT_OP_KERNEL_ARG_DEFAULT_TYPE_LIST(
@ -66,11 +80,16 @@ using Pad11Types = ORT_OP_KERNEL_ARG_DEFAULT_TYPE_LIST(
kCpuExecutionProvider, kOnnxDomain, Pad, 11, Input, 0);
using EnabledPad11Types = ORT_OP_KERNEL_ARG_ENABLED_TYPE_LIST(
kCpuExecutionProvider, kOnnxDomain, Pad, 11, Input, 0);
using Pad13Types = ORT_OP_KERNEL_ARG_DEFAULT_TYPE_LIST(
kCpuExecutionProvider, kOnnxDomain, Pad, 13, Input, 0);
using EnabledPad13Types = ORT_OP_KERNEL_ARG_ENABLED_TYPE_LIST(
kCpuExecutionProvider, kOnnxDomain, Pad, 13, Input, 0);
using AllEnabledPadTypes =
utils::TypeSetUnion<
EnabledPad2Types,
EnabledPad11Types>;
EnabledPad11Types,
EnabledPad13Types>;
// only float type is supported for opset-10
ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
@ -98,10 +117,11 @@ ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
ONNX_CPU_OPERATOR_KERNEL(
Pad,
13,
KernelDefBuilder().TypeConstraint(
"T",
BuildKernelDefConstraintsFromTypeList<Pad11Types>(),
BuildKernelDefConstraintsFromTypeList<EnabledPad11Types>()),
KernelDefBuilder()
.TypeConstraint(
"T",
BuildKernelDefConstraintsFromTypeList<Pad13Types>(),
BuildKernelDefConstraintsFromTypeList<EnabledPad13Types>()),
Pad);
// This is the general padding method to n-dimensionally do edge or reflection padding (based on the inputDelta values)

View file

@ -166,13 +166,26 @@ Status Transpose::DoTranspose(const cudaDeviceProp& prop,
if (CanDoTranspose3D(new_rank, new_input_dims, new_permutations)) {
return Transpose3DImpl(stream, element_size, input_shape, tmp_input_strides,
input.DataRaw(), output.MutableDataRaw(), output.Shape().Size());
} else if (CanDoTranspose4D(prop, element_size, new_rank, new_input_dims, new_permutations)) {
} else if (CanDoTranspose4DParallelizeMultipleElementsPerThreadInInnermostDim(
prop, element_size, new_rank, new_input_dims, new_permutations)) {
TArray<int64_t> tmp_output_strides(new_rank);
for (auto i = 0; i < new_rank; i++) {
tmp_output_strides[i] = new_output_strides[new_permutations[i]];
}
return Transpose4DImpl(stream, element_size, input_shape, tmp_input_strides, input.DataRaw(),
tmp_output_strides, output.MutableDataRaw(), gsl::narrow<int>(output.Shape().Size()));
return Transpose4DParallelizeMultipleElementsPerThreadInInnermostDim(
stream, element_size, input_shape, tmp_input_strides, input.DataRaw(),
tmp_output_strides, output.MutableDataRaw(), gsl::narrow<int>(output.Shape().Size()));
} else if (CanDoTranspose4DParallelizeOneElementPerThread(
prop, element_size, new_rank, new_input_dims, new_permutations)) {
// Trying to see if we can still do (best effort) more optimized transposing
// for the 4-D case before falling back to the generic case
TArray<int64_t> tmp_output_strides(new_rank);
for (auto i = 0; i < new_rank; i++) {
tmp_output_strides[i] = new_output_strides[new_permutations[i]];
}
return Transpose4DParallelizeOneElementPerThread(
stream, element_size, input_shape, tmp_input_strides, input.DataRaw(),
tmp_output_strides, output.MutableDataRaw(), gsl::narrow<int>(output.Shape().Size()));
}
// General cases

View file

@ -80,9 +80,10 @@ Status Transpose3DImpl(cudaStream_t stream, size_t element_size,
}
template <int element_size>
__global__ void Transpose4DKernel(const TArray<int64_t> input_strides, const void* input_data,
const TArray<int64_t> output_strides, void* output_data,
CUDA_LONG N) {
__global__ void Transpose4DKernelParallelizeMultipleElementsPerThreadInInnermostDim(
const TArray<int64_t> input_strides, const void* input_data,
const TArray<int64_t> output_strides, void* output_data,
CUDA_LONG N) {
// output coordinates will be: blockIdx.y, blockIdx.x, threadIdx.y, threadIdx.x
CUDA_LONG input_index = (blockIdx.y * input_strides[0] +
blockIdx.x * input_strides[1] +
@ -104,59 +105,69 @@ __global__ void Transpose4DKernel(const TArray<int64_t> input_strides, const voi
}
}
bool CanDoTranspose4D(const cudaDeviceProp& prop,
size_t element_size,
int32_t rank,
const std::vector<int64_t>& input_dims,
const std::vector<size_t>& permutations) {
bool CanDoTranspose4DParallelizeMultipleElementsPerThreadInInnermostDim(const cudaDeviceProp& prop,
size_t element_size,
int32_t rank,
const std::vector<int64_t>& input_dims,
const std::vector<size_t>& permutations) {
if (rank == 4 &&
// the permutations is not on the last dimension.
permutations[rank - 1] == (rank - 1)) {
// The block size will be set based on the last two dimensions of 4D tensor.
permutations[3] == 3) {
// The block size will be set based on the outer-most two dimensions of 4D tensor.
// the number threads per block will be calculated as below.
unsigned int num_elements_per_thread = 4 * sizeof(int) / static_cast<unsigned int>(element_size); // int4 is used in the kernel to access data.
int64_t num_elements_in_last_two_dimensions = input_dims[rank - 2] * input_dims[rank - 1];
int64_t num_elements_in_last_two_dimensions = input_dims[2] * input_dims[3];
int64_t num_threads_per_block = num_elements_in_last_two_dimensions / num_elements_per_thread;
if (((num_elements_in_last_two_dimensions & (num_elements_per_thread - 1)) == 0) &&
num_threads_per_block <= prop.maxThreadsPerBlock &&
num_threads_per_block >= prop.warpSize &&
// num_threads_per_block must be aligned with warp size: 32
((num_threads_per_block & (prop.warpSize - 1)) == 0)) {
// num_threads_per_block must be a multiple of warp size (32)
((num_threads_per_block & (prop.warpSize - 1)) == 0) &&
// input_dims[3] must be a multiple of `num_elements_per_thread`
((input_dims[3] % num_elements_per_thread) == 0)) {
return true;
}
}
return false;
}
Status Transpose4DImpl(cudaStream_t stream, size_t element_size, const TArray<int64_t>& input_shape, const TArray<int64_t>& input_strides, const void* input_data,
const TArray<int64_t>& output_strides, void* output_data, int N) {
Status Transpose4DParallelizeMultipleElementsPerThreadInInnermostDim(
cudaStream_t stream, size_t element_size,
const TArray<int64_t>& input_shape, const TArray<int64_t>& input_strides,
const void* input_data, const TArray<int64_t>& output_strides,
void* output_data, int N) {
unsigned int num_elements_per_thread = 4 * sizeof(int) / static_cast<unsigned int>(element_size); // int4 is used in the kernel to access data.
dim3 block_size(static_cast<unsigned int>(input_shape[3] / num_elements_per_thread), static_cast<unsigned int>(input_shape[2]));
dim3 grid_size(static_cast<unsigned int>(input_shape[1]), static_cast<unsigned int>(input_shape[0]));
switch (element_size) {
case sizeof(int8_t):
Transpose4DKernel<sizeof(int8_t)><<<grid_size, block_size, 0, stream>>>(
input_strides, input_data,
output_strides, output_data, N / num_elements_per_thread);
Transpose4DKernelParallelizeMultipleElementsPerThreadInInnermostDim<sizeof(int8_t)>
<<<grid_size, block_size, 0, stream>>>(
input_strides, input_data,
output_strides, output_data, N / num_elements_per_thread);
break;
case sizeof(int16_t):
Transpose4DKernel<sizeof(int16_t)><<<grid_size, block_size, 0, stream>>>(
input_strides, input_data,
output_strides, output_data, N / num_elements_per_thread);
Transpose4DKernelParallelizeMultipleElementsPerThreadInInnermostDim<sizeof(int16_t)>
<<<grid_size, block_size, 0, stream>>>(
input_strides, input_data,
output_strides, output_data, N / num_elements_per_thread);
break;
case sizeof(int32_t):
Transpose4DKernel<sizeof(int32_t)><<<grid_size, block_size, 0, stream>>>(
input_strides, input_data,
output_strides, output_data, N / num_elements_per_thread);
Transpose4DKernelParallelizeMultipleElementsPerThreadInInnermostDim<sizeof(int32_t)>
<<<grid_size, block_size, 0, stream>>>(
input_strides, input_data,
output_strides, output_data, N / num_elements_per_thread);
break;
case sizeof(int64_t):
Transpose4DKernel<sizeof(int64_t)><<<grid_size, block_size, 0, stream>>>(
input_strides, input_data,
output_strides, output_data, N / num_elements_per_thread);
Transpose4DKernelParallelizeMultipleElementsPerThreadInInnermostDim<sizeof(int64_t)>
<<<grid_size, block_size, 0, stream>>>(
input_strides, input_data,
output_strides, output_data, N / num_elements_per_thread);
break;
default:
// User will not hit this as this kernel is for fixed element size tensors only
return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "Type not supported for transpose on CUDA. Element size was ",
element_size);
}
@ -164,6 +175,77 @@ Status Transpose4DImpl(cudaStream_t stream, size_t element_size, const TArray<in
return Status::OK();
}
__global__ void Transpose4DKernelParallelizeOneElementPerThread(
const TArray<int64_t> input_strides, const int8_t* input_data,
const TArray<int64_t> output_strides, int8_t* output_data,
size_t element_size,
CUDA_LONG N) {
CUDA_LONG input_index = blockIdx.y * input_strides[0] +
blockIdx.x * input_strides[1] +
threadIdx.y * input_strides[2] +
threadIdx.x * input_strides[3];
CUDA_LONG output_index = blockIdx.y * output_strides[0] +
blockIdx.x * output_strides[1] +
threadIdx.y * output_strides[2] +
threadIdx.x * output_strides[3];
if (input_index < N && output_index < N) {
const int8_t* input_data_to_be_copied = input_data + (input_index * element_size);
int8_t* output_data_to_be_copied = output_data + (output_index * element_size);
// copy over the bytes
for (size_t iter = 0; iter < element_size; ++iter) {
*output_data_to_be_copied++ = *input_data_to_be_copied++;
}
}
}
bool CanDoTranspose4DParallelizeOneElementPerThread(const cudaDeviceProp& prop,
size_t element_size,
int32_t rank,
const std::vector<int64_t>& input_dims,
const std::vector<size_t>& permutations) {
if (rank == 4) {
// The block size will be set based on the outer-most two dimensions of 4D tensor.
// the number threads per block will be calculated as below.
int64_t number_of_threads_per_block = input_dims[2] * input_dims[3];
if (number_of_threads_per_block <= prop.maxThreadsPerBlock &&
number_of_threads_per_block >= prop.warpSize &&
// num_threads_per_block must be a multiple of warp size (32)
((number_of_threads_per_block & (prop.warpSize - 1)) == 0)) {
return true;
}
}
return false;
}
Status Transpose4DParallelizeOneElementPerThread(
cudaStream_t stream, size_t element_size,
const TArray<int64_t>& input_shape, const TArray<int64_t>& input_strides,
const void* input_data, const TArray<int64_t>& output_strides,
void* output_data, int N) {
if (element_size != sizeof(int8_t) &&
element_size != sizeof(int16_t) &&
element_size != sizeof(int32_t) &&
element_size != sizeof(int64_t)) {
// User will not hit this as this kernel is for fixed element size tensors only
return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "Type not supported for transpose on CUDA. Element size was ",
element_size);
}
dim3 block_size(static_cast<unsigned int>(input_shape[3]), static_cast<unsigned int>(input_shape[2]));
dim3 grid_size(static_cast<unsigned int>(input_shape[1]), static_cast<unsigned int>(input_shape[0]));
Transpose4DKernelParallelizeOneElementPerThread<<<grid_size, block_size, 0, stream>>>(
input_strides, reinterpret_cast<const int8_t*>(input_data),
output_strides, reinterpret_cast<int8_t*>(output_data),
element_size, N);
return Status::OK();
}
template <typename T>
__global__ void TransposeKernel(int32_t shape_rank, const TArray<int64_t> input_strides,
const T* input_data, const TArray<fast_divmod> output_strides, T* output_data, CUDA_LONG N) {

View file

@ -11,13 +11,25 @@ namespace cuda {
bool CanDoTranspose3D(int32_t rank, const std::vector<int64_t>& input_dims, const std::vector<size_t>& permutations);
Status Transpose3DImpl(cudaStream_t stream, size_t element_size, const TArray<int64_t>& input_shape, const TArray<int64_t>& input_strides, const void* input_data,
void* output_data, int64_t N);
bool CanDoTranspose4D(const cudaDeviceProp& prop,
size_t element_size,
int32_t rank,
const std::vector<int64_t>& input_dims,
const std::vector<size_t>& permutations);
Status Transpose4DImpl(cudaStream_t stream, size_t element_size, const TArray<int64_t>& input_shape, const TArray<int64_t>& input_strides, const void* input_data,
const TArray<int64_t>& output_strides, void* output_data, int N);
bool CanDoTranspose4DParallelizeMultipleElementsPerThreadInInnermostDim(const cudaDeviceProp& prop,
size_t element_size,
int32_t rank,
const std::vector<int64_t>& input_dims,
const std::vector<size_t>& permutations);
Status Transpose4DParallelizeMultipleElementsPerThreadInInnermostDim(cudaStream_t stream, size_t element_size, const TArray<int64_t>& input_shape,
const TArray<int64_t>& input_strides, const void* input_data,
const TArray<int64_t>& output_strides, void* output_data, int N);
bool CanDoTranspose4DParallelizeOneElementPerThread(const cudaDeviceProp& prop,
size_t element_size,
int32_t rank,
const std::vector<int64_t>& input_dims,
const std::vector<size_t>& permutations);
Status Transpose4DParallelizeOneElementPerThread(cudaStream_t stream, size_t element_size, const TArray<int64_t>& input_shape,
const TArray<int64_t>& input_strides, const void* input_data,
const TArray<int64_t>& output_strides, void* output_data, int N);
Status TransposeImpl(cudaStream_t stream, size_t element_size, int32_t shape_rank, const TArray<int64_t>& input_strides,
const void* input_data, const TArray<fast_divmod>& fdm_output_strides, void* output_data, int N);
} // namespace cuda

View file

@ -62,16 +62,16 @@ Status TransposeWithRocblas(hipStream_t stream, rocblas_handle rocblas_handle, c
HipT* output_data = reinterpret_cast<HipT*>(output.MutableData<T>());
ROCBLAS_RETURN_IF_ERROR(
rocblasTransposeHelper(stream,
rocblas_handle,
rocblas_operation_transpose, rocblas_operation_transpose, M, N,
&one,
input_data,
N,
&zero,
input_data,
N,
output_data,
M));
rocblas_handle,
rocblas_operation_transpose, rocblas_operation_transpose, M, N,
&one,
input_data,
N,
&zero,
input_data,
N,
output_data,
M));
return Status::OK();
}
@ -128,25 +128,25 @@ Status Transpose::DoTranspose(const hipDeviceProp_t& prop,
new_permutations[j] -= 1;
}
}
for (auto j = i+1; j < new_rank; j++) {
new_permutations[j-1] = new_permutations[j];
for (auto j = i + 1; j < new_rank; j++) {
new_permutations[j - 1] = new_permutations[j];
}
// update input dims
new_input_dims[prev] *= new_input_dims[curr];
new_input_dims[curr] = 1;
for (auto j = static_cast<int32_t>(curr+1); j < new_rank; j++) {
new_input_dims[j-1] = new_input_dims[j];
for (auto j = static_cast<int32_t>(curr + 1); j < new_rank; j++) {
new_input_dims[j - 1] = new_input_dims[j];
}
new_input_dims[new_rank-1] = 1;
new_input_dims[new_rank - 1] = 1;
// update output dims
new_output_dims[i-1] *= new_output_dims[i];
new_output_dims[i - 1] *= new_output_dims[i];
new_output_dims[i] = 1;
for (auto j = i+1; j < new_rank; j++) {
new_output_dims[j-1] = new_output_dims[j];
for (auto j = i + 1; j < new_rank; j++) {
new_output_dims[j - 1] = new_output_dims[j];
}
new_output_dims[new_rank-1] = 1;
new_output_dims[new_rank - 1] = 1;
new_rank--;
}
@ -166,13 +166,26 @@ Status Transpose::DoTranspose(const hipDeviceProp_t& prop,
if (CanDoTranspose3D(new_rank, new_input_dims, new_permutations)) {
return Transpose3DImpl(stream, element_size, input_shape, tmp_input_strides,
input.DataRaw(), output.MutableDataRaw(), output.Shape().Size());
} else if (CanDoTranspose4D(prop, element_size, new_rank, new_input_dims, new_permutations)) {
} else if (CanDoTranspose4DParallelizeMultipleElementsPerThreadInInnermostDim(
prop, element_size, new_rank, new_input_dims, new_permutations)) {
TArray<int64_t> tmp_output_strides(new_rank);
for (auto i = 0; i < new_rank; i++) {
tmp_output_strides[i] = new_output_strides[new_permutations[i]];
}
return Transpose4DImpl(stream, element_size, input_shape, tmp_input_strides, input.DataRaw(),
tmp_output_strides, output.MutableDataRaw(), output.Shape().Size());
return Transpose4DParallelizeMultipleElementsPerThreadInInnermostDim(
stream, element_size, input_shape, tmp_input_strides, input.DataRaw(),
tmp_output_strides, output.MutableDataRaw(), gsl::narrow<int>(output.Shape().Size()));
} else if (CanDoTranspose4DParallelizeOneElementPerThread(
prop, element_size, new_rank, new_input_dims, new_permutations)) {
// Trying to see if we can still do (best effort) more optimized transposing
// for the 4-D case before falling back to the generic case
TArray<int64_t> tmp_output_strides(new_rank);
for (auto i = 0; i < new_rank; i++) {
tmp_output_strides[i] = new_output_strides[new_permutations[i]];
}
return Transpose4DParallelizeOneElementPerThread(
stream, element_size, input_shape, tmp_input_strides, input.DataRaw(),
tmp_output_strides, output.MutableDataRaw(), gsl::narrow<int>(output.Shape().Size()));
}
// General cases

View file

@ -7,7 +7,6 @@ from pathlib import Path
class ONNXModel:
def __init__(self, model):
self.model = model
self.node_name_counter = {}
def nodes(self):
return self.model.graph.node
@ -260,38 +259,49 @@ class ONNXModel:
def topological_sort(self):
deps_count = [0]*len(self.nodes()) # dependency count of each node
deps_to_nodes = {} # input to node indice
sorted_nodes = [] # initialize sorted_nodes
for node_idx, node in enumerate(self.nodes()):
# CANNOT use len(node.input) directly because input can be optional
deps_count[node_idx] = sum(1 for _ in node.input if _ )
if deps_count[node_idx] == 0: # Constant doesn't depend on any inputs
sorted_nodes.append(self.nodes()[node_idx])
continue
for input_name in node.input:
if input_name not in deps_to_nodes:
deps_to_nodes[input_name] = [node_idx]
else:
deps_to_nodes[input_name].append(node_idx)
# initialize sorted_nodes
sorted_nodes = []
for input in itertools.chain(self.initializer(), self.model.graph.input):
if input.name in deps_to_nodes:
for node_idx in deps_to_nodes[input.name]:
initializer_names = [init.name for init in self.initializer()]
graph_input_names = [input.name for input in self.model.graph.input]
input_names = initializer_names + graph_input_names
input_names.sort()
prev_input_name = None
for input_name in input_names:
if prev_input_name == input_name:
continue
prev_input_name = input_name
if input_name in deps_to_nodes:
for node_idx in deps_to_nodes[input_name]:
deps_count[node_idx] = deps_count[node_idx] - 1
if deps_count[node_idx] == 0:
sorted_nodes.append(self.nodes()[node_idx])
s = 0
e = len(sorted_nodes)
start = 0
end = len(sorted_nodes)
while s < e:
for output in sorted_nodes[s].output:
while start < end:
for output in sorted_nodes[start].output:
if output in deps_to_nodes:
for node_idx in deps_to_nodes[output]:
deps_count[node_idx] = deps_count[node_idx] - 1
if deps_count[node_idx] == 0:
sorted_nodes.append(self.nodes()[node_idx])
e = e + 1
s = s + 1
end = end + 1
start = start + 1
assert(e == len(self.graph().node)), "Graph is not a DAG"
assert(end == len(self.graph().node)), "Graph is not a DAG"
self.graph().ClearField('node')
self.graph().node.extend(sorted_nodes)
self.graph().node.extend(sorted_nodes)

View file

@ -112,7 +112,7 @@ class FusionLayerNormalization(Fusion):
inputs=[node.input[0], weight_input, bias_input],
outputs=[last_add_node.output[0]],
name=self.model.create_node_name("LayerNormalization",
name_prefix="SkipLayerNorm"))
name_prefix="LayerNorm"))
normalize_node.attribute.extend([helper.make_attribute("epsilon", float(add_weight))])
self.nodes_to_add.append(normalize_node)
self.node_name_to_graph_name[normalize_node.name] = self.this_graph_name

View file

@ -3,7 +3,7 @@
# Licensed under the MIT License.
#--------------------------------------------------------------------------
from typing import List, Tuple
from typing import List, Tuple, Dict
import logging
import os
import sys
@ -19,7 +19,7 @@ logger = logging.getLogger(__name__)
class OnnxModel:
def __init__(self, model):
self.model = model
self.node_name_counter = {}
self._node_name_suffix: Dict[str, int] = {} # key is node name prefix, value is the last suffix generated
self.shape_infer_helper = None
self.all_graphs = None
@ -553,25 +553,39 @@ class OnnxModel:
cast_node.attribute.extend([helper.make_attribute("to", int(TensorProto.FLOAT))])
self.add_node(cast_node)
# create a new name for node
def create_node_name(self, op_type, name_prefix=None):
if op_type in self.node_name_counter:
self.node_name_counter[op_type] += 1
"""Create a unique node name that starts with a prefix (default is operator type).
The name will not be duplicated with any name that generated or existed in current graphs.
Args:
op_type (str): operator type
name_prefix (str, optional): prefix of node name. Defaults to None.
Returns:
str: node name
"""
if name_prefix:
prefix = name_prefix if name_prefix.endswith("_") else (name_prefix + "_")
else:
self.node_name_counter[op_type] = 1
prefix = op_type + "_"
if name_prefix is not None:
full_name = name_prefix + str(self.node_name_counter[op_type])
suffix: int = 0
if prefix in self._node_name_suffix:
suffix = self._node_name_suffix[prefix] + 1
else:
full_name = op_type + "_" + str(self.node_name_counter[op_type])
# Check existed node name only once for a prefix as we assume create_node_name is called for every new node in fusion.
for node in self.nodes():
if node.name and node.name.startswith(prefix):
try:
index = int(node.name[len(prefix):])
suffix = max(index + 1, suffix)
except ValueError:
continue
# Check whether the name is taken:
nodes = self.get_nodes_by_op_type(op_type)
for node in nodes:
if node.name == full_name:
raise Exception("Node name already taken:", full_name)
# Record the generated suffix so that we can avoid generating duplicated name.
self._node_name_suffix[prefix] = suffix
return full_name
return prefix + str(suffix)
def find_graph_input(self, input_name):
for input in self.model.graph.input:

View file

@ -26,8 +26,7 @@ static void RunOnnxOpsetTypedTest(
if (opset >= 11) {
test.AddInput<int64_t>("pads", {static_cast<int64_t>(pads.size())}, pads);
test.AddInput<T>("value", {1}, {value});
}
else {
} else {
test.AddAttribute("pads", pads);
test.AddAttribute("value", static_cast<float>(value));
}
@ -35,8 +34,7 @@ static void RunOnnxOpsetTypedTest(
if (opset >= 11) {
// TensorRT do not yet support opset-11 and builds break on this test, hence exclude the EP
test.Run(expect, error_msg, {kTensorrtExecutionProvider});
}
else {
} else {
#if defined(OPENVINO_CONFIG_MYRIAD) || defined(OPENVINO_CONFIG_VAD_M)
test.Run(expect, error_msg, {kOpenVINOExecutionProvider});
#else
@ -56,7 +54,7 @@ static void RunAllOpsetAllDomainPadTests(
std::string mode = "constant",
OpTester::ExpectResult expect = OpTester::ExpectResult::kExpectSuccess,
const std::string& error_msg = "") {
// ONNX domain opset-11 is the only one to support all data types
// Test opset-11 and opset-13 kernels of Pad
RunOnnxOpsetTypedTest<T, 11>(input_dims,
input,
pads,
@ -64,8 +62,16 @@ static void RunAllOpsetAllDomainPadTests(
output_dims,
output,
mode, expect, error_msg);
RunOnnxOpsetTypedTest<T, 13>(input_dims,
input,
pads,
value,
output_dims,
output,
mode, expect, error_msg);
}
template<>
template <>
void RunAllOpsetAllDomainPadTests<>(
const std::vector<int64_t>& input_dims,
const std::vector<double>& input,
@ -76,7 +82,7 @@ void RunAllOpsetAllDomainPadTests<>(
std::string mode,
OpTester::ExpectResult expect,
const std::string& error_msg) {
// ONNX domain supports double type
// Test opset-10, opset-11 and opset-13 kernels of Pad (for double type)
RunOnnxOpsetTypedTest<double, 10>(input_dims,
input,
pads,
@ -84,6 +90,7 @@ void RunAllOpsetAllDomainPadTests<>(
output_dims,
output,
mode, expect, error_msg);
RunOnnxOpsetTypedTest<double, 11>(input_dims,
input,
pads,
@ -91,9 +98,18 @@ void RunAllOpsetAllDomainPadTests<>(
output_dims,
output,
mode, expect, error_msg);
RunOnnxOpsetTypedTest<double, 13>(input_dims,
input,
pads,
value,
output_dims,
output,
mode, expect, error_msg);
}
// There is only support for float type for MSDomain kernel in ORT
template<>
template <>
void RunAllOpsetAllDomainPadTests<>(
const std::vector<int64_t>& input_dims,
const std::vector<float>& input,
@ -104,6 +120,7 @@ void RunAllOpsetAllDomainPadTests<>(
std::string mode,
OpTester::ExpectResult expect,
const std::string& error_msg) {
// Test opset-10, opset-11 and opset-13 kernels of Pad (for float type)
RunOnnxOpsetTypedTest<float, 10>(input_dims,
input,
pads,
@ -111,6 +128,7 @@ void RunAllOpsetAllDomainPadTests<>(
output_dims,
output,
mode, expect, error_msg);
RunOnnxOpsetTypedTest<float, 11>(input_dims,
input,
pads,
@ -119,6 +137,14 @@ void RunAllOpsetAllDomainPadTests<>(
output,
mode, expect, error_msg);
RunOnnxOpsetTypedTest<float, 13>(input_dims,
input,
pads,
value,
output_dims,
output,
mode, expect, error_msg);
#ifndef DISABLE_CONTRIB_OPS
// MSFT domain opset-1 (contrib op)
@ -129,7 +155,7 @@ void RunAllOpsetAllDomainPadTests<>(
test3.AddInput<float>("value", {1}, {value});
test3.AddOutput<float>("output", output_dims, output);
//TensorRT does not support pads as an input
test3.Run(expect, error_msg, {kTensorrtExecutionProvider,kOpenVINOExecutionProvider});
test3.Run(expect, error_msg, {kTensorrtExecutionProvider, kOpenVINOExecutionProvider});
#endif
}
@ -679,19 +705,19 @@ TYPED_TEST(PadOpTest, Pad_Constant_DimWithZeroInput) {
{T(1), T(1), T(1), T(1), T(1), T(1), T(1), T(1)});
}
// Added output shape verification b/w the output shape generated by operator specific ONNX inference and
// the output shape generated by operator specific ORT implementation. After adding this verification,
// the output shape generated by operator specific ORT implementation. After adding this verification,
// this test logs warning as validation fails for 2 data types out of 8 data types i.e. Float and Double.
// Reason:
// Pad ORT implementation output shape does not match with Pad ONNX inference function output shape.
//
// For Float and Double this test gets executed for 2 different opset version, 10 and 11. Specifically this
// test is failing for opset version 10.
// Investigation Analysis: Different ONNX inference class/method gets executed per opset version. Main difference b/w the 2
//
// For Float and Double this test gets executed for 2 different opset version, 10 and 11. Specifically this
// test is failing for opset version 10.
// Investigation Analysis: Different ONNX inference class/method gets executed per opset version. Main difference b/w the 2
// pad operator ONNX inference class/method is:
// Older Pad operator ONNX inference: Accepts "pads and values" as attribute.
// Newer Pad operator ONNX inference: Accetps "pads and values" as input.
// For newer version, "pads & values" fields have not been added as initializer, thus instead of shape
// inference, rank inference gets triggered. Whereas, in older version shape inference gets executed
// For newer version, "pads & values" fields have not been added as initializer, thus instead of shape
// inference, rank inference gets triggered. Whereas, in older version shape inference gets executed
// as "pads & values" fields have been added as attribute.
// In order to remove the warning, shape inference methods needs to be fixed.
@ -743,5 +769,15 @@ TYPED_TEST(PadOpTest, Pad_Reflect_DimWithZeroInput) {
"Cannot use 'reflect' mode to pad dimension with a value of 0. Input shape:{0,2,1}");
}
TEST(PadOpTest, BoolType) {
OpTester test("Pad", 13);
test.AddAttribute("mode", "constant");
test.AddInput<bool>("data", {3, 2}, {true, false, true, false, true, false});
test.AddInput<int64_t>("pads", {4}, {0, 2, 0, 0});
test.AddInput<bool>("value", {1}, {true});
test.AddOutput<bool>("output", {3, 4}, {true, true, true, false, true, true, true, false, true, true, true, false});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});
}
} // namespace test
} // namespace onnxruntime

View file

@ -590,26 +590,34 @@ static void TestTranspose(
test.CompareWithCPU(kGpuExecutionProvider, error_tolerance);
}
TEST(TransposeOpTest, Transpose0213) {
TEST(TransposeOpTest, Transpose0213) { // Will trigger Transpose4DParallelizeMultipleElementsPerThreadInInnermostDim()
const std::vector<int64_t> X_dims{64, 128, 16, 64};
const std::vector<int64_t> perm{0, 2, 1, 3};
const std::vector<int64_t> Y_dims{64, 16, 128, 64};
TestTranspose(perm, X_dims, Y_dims);
}
TEST(TransposeOpTest, Transpose0231) {
TEST(TransposeOpTest, Transpose0213_V2) { // Will trigger Transpose4DParallelizeOneElementPerThread()
const std::vector<int64_t> X_dims{64, 128, 64, 2};
const std::vector<int64_t> perm{0, 2, 1, 3};
const std::vector<int64_t> Y_dims{64, 64, 128, 2};
TestTranspose(perm, X_dims, Y_dims);
}
TEST(TransposeOpTest, Transpose0231) { // Will trigger Transpose3DImpl() because of "flattening" of dims 2 and 3 into one dim
const std::vector<int64_t> X_dims{64, 128, 16, 64};
const std::vector<int64_t> perm{0, 2, 3, 1};
const std::vector<int64_t> Y_dims{64, 16, 64, 128};
TestTranspose(perm, X_dims, Y_dims);
}
TEST(TransposeOpTest, Transpose0312) {
TEST(TransposeOpTest, Transpose0312) { // Will trigger Transpose3DImpl() because of "flattening" of dims 1 and 2 into one dim
const std::vector<int64_t> X_dims{64, 16, 64, 128};
const std::vector<int64_t> perm{0, 3, 1, 2};
const std::vector<int64_t> Y_dims{64, 128, 16, 64};
TestTranspose(perm, X_dims, Y_dims);
}
#endif
} // namespace test

View file

@ -65,6 +65,28 @@ class TestONNXModel(unittest.TestCase):
model = helper.make_model(graph, opset_imports=[helper.make_opsetid("", 13)])
onnx.save(model, model_path)
def construct_model_Constant(self, model_path):
# (input) Constant
# \ /
# \ /
# \ /
# \ /
# Add
# |
# (output)
initializers = []
input = helper.make_tensor_value_info('input', TensorProto.FLOAT, [4, 8, 12])
output = helper.make_tensor_value_info('output', TensorProto.FLOAT, [4, 8, 12])
# make nodes
constant_node = onnx.helper.make_node('Constant', [], ['const_output'], value_float=42.0)
add_node = onnx.helper.make_node('Add', ['input', 'const_output'], ['output'], name='Add')
graph = helper.make_graph([add_node, constant_node],
'onnx_model_test', [input], [output], initializer=initializers)
model = helper.make_model(graph, opset_imports=[helper.make_opsetid("", 13)])
onnx.save(model, model_path)
def test_topo_sort(self):
test_model_path = 'onnx_model_topo_sort.onnx'
self.construct_model(test_model_path)
@ -73,5 +95,13 @@ class TestONNXModel(unittest.TestCase):
onnx_model.topological_sort()
check_op_type_order(self, onnx_model.model, ['GRU', 'Conv', 'Conv', 'Relu', 'Add'])
def test_topo_sort_constant(self):
test_model_path = 'onnx_model_topo_sort_constant.onnx'
self.construct_model_Constant(test_model_path)
onnx_model = ONNXModel(onnx.load(test_model_path))
check_op_type_order(self, onnx_model.model, ['Add', 'Constant'])
onnx_model.topological_sort()
check_op_type_order(self, onnx_model.model, ['Constant', 'Add'])
if __name__ == '__main__':
unittest.main()

View file

@ -1461,7 +1461,7 @@
],
[
"Pad ai.onnx CPUExecutionProvider",
9596174091174553032
12904240253005862936
],
[
"Pad ai.onnx CPUExecutionProvider",

View file

@ -11,7 +11,29 @@ from packaging import version
################################################################################
ONNX_OPSET_VERSION = 12
MINIMUM_TORCH_VERSION_STR = '1.8.1'
TORCH_CPP_BUILD_DIR = os.path.join(os.path.dirname(__file__),'torch_inline_extensions')
# Use one of the available directories as Torch CPP extension in the following order:
# 1) Path at listed at TORCH_EXTENSIONS_DIR environment variable
# 2) Default Python package dir
# 3) <Home directory>/.cache
home_dir = os.path.expanduser("~")
python_package_dir = os.path.dirname(__file__)
torch_extensions_dir = os.environ.get('TORCH_EXTENSIONS_DIR')
TORCH_CPP_BUILD_DIR = os.path.join(python_package_dir,'torch_inline_extensions')
TORCH_CPP_BUILD_DIR_BACKUP = os.path.join(home_dir, '.cache', 'torch_ort_extensions')
if torch_extensions_dir is not None and os.access(torch_extensions_dir, os.X_OK | os.W_OK):
TORCH_CPP_BUILD_DIR = torch_extensions_dir
elif not os.access(python_package_dir, os.X_OK | os.W_OK):
if os.access(home_dir, os.X_OK | os.W_OK):
TORCH_CPP_BUILD_DIR = TORCH_CPP_BUILD_DIR_BACKUP
else:
extra_message = ''
if torch_extensions_dir:
extra_message = 'or the path pointed by the TORCH_EXTENSIONS_DIR environment variable '
raise PermissionError('ORTModule could not find a writable directory to cache its internal files.',
f'Make {python_package_dir} or {home_dir} {extra_message}writable and try again.')
# Check whether Torch C++ extension compilation was aborted in previous runs
if not os.path.exists(TORCH_CPP_BUILD_DIR):
@ -19,19 +41,19 @@ if not os.path.exists(TORCH_CPP_BUILD_DIR):
elif os.path.exists(os.path.join(TORCH_CPP_BUILD_DIR,'lock')):
print("WARNING: ORTModule detected PyTorch CPP extension's lock file during initialization, "
"which can cause unexpected hangs. "
f"Delete {os.path.join(TORCH_CPP_BUILD_DIR,'lock')} to prevent unexpected behavior.")
f"Delete {os.path.join(TORCH_CPP_BUILD_DIR,'lock')} if a hang occurs.")
# Verify proper PyTorch is installed before proceding to ONNX Runtime initializetion
# Verify proper PyTorch is installed before proceding to ONNX Runtime initialization
try:
import torch
torch_version = version.parse(torch.__version__.split('+')[0])
minimum_torch_version = version.parse(MINIMUM_TORCH_VERSION_STR)
if torch_version < minimum_torch_version:
raise RuntimeError(
f'ONNXRuntime ORTModule frontend requires PyTorch version greater or equal to {MINIMUM_TORCH_VERSION_STR}, '
f'ONNX Runtime ORTModule frontend requires PyTorch version greater or equal to {MINIMUM_TORCH_VERSION_STR}, '
f'but version {torch.__version__} was found instead.')
except:
raise(f'PyTorch {MINIMUM_TORCH_VERSION_STR} must be installed in order to run ONNXRuntime ORTModule frontend!')
raise(f'PyTorch {MINIMUM_TORCH_VERSION_STR} must be installed in order to run ONNX Runtime ORTModule frontend!')
# ORTModule must be loaded only after all validation passes
from .ortmodule import ORTModule

View file

@ -97,3 +97,10 @@ def _create_iobinding(io_binding, inputs, model, device):
for value_info in model.graph.output:
io_binding.bind_output(value_info.name, device.type, device_id=get_device_index(device))
class _PytorchModuleMetadata():
"""Encapsulates modules and allows easy access as required"""
def __init__(self, original_module, flattened_module):
self.original_module = original_module
self.flattened_module = flattened_module

View file

@ -5,12 +5,13 @@
from . import _io
from ._graph_execution_manager_factory import GraphExecutionManagerFactory
from ._utils import _PytorchModuleMetadata
from onnxruntime.training import register_custom_ops_pytorch_exporter
import functools
import torch
from typing import Iterator, Optional, Tuple, TypeVar
from typing import Iterator, Optional, Tuple, TypeVar, Set, Callable
# Needed to override PyTorch methods
T = TypeVar('T', bound='Module')
@ -51,12 +52,11 @@ class ORTModule(torch.nn.Module):
register_custom_ops_pytorch_exporter.register_custom_op(is_ortmodule=True)
# User module is wrapped to use its initializers and save computed gradients
self._original_module = module
# along with the module that flattens both input and output of the user module
# inside _PytorchModuleMetadata
self._module_metadata = _PytorchModuleMetadata(module, _io._FlattenedModule(module))
# Get the module that flattens both input and output
self._flattened_module = _io._FlattenedModule(self._original_module)
self._execution_manager = GraphExecutionManagerFactory(self._flattened_module)
self._execution_manager = GraphExecutionManagerFactory(self._module_metadata.flattened_module)
# IMPORTANT: DO NOT add code here
# This declaration is for automatic document generation purposes only
@ -65,57 +65,82 @@ class ORTModule(torch.nn.Module):
'''Dummy documentation for forward method'''
...
def _apply(self, fn):
"""Override original method to delegate execution to the flattened PyTorch user module"""
# Delegation must happen to _flattened_module since methods depend on
# _apply to recursively apply the internal setting changes
self._module_metadata.flattened_module._apply(fn)
return self
def apply(self: T, fn: Callable[['Module'], None]) -> T:
"""Override original method to delegate execution to the flattened PyTorch user module"""
# Delegation must happen to _flattened_module since methods depend on
# apply to recursively apply the internal setting changes
self._module_metadata.flattened_module.apply(fn)
return self
def _is_training(self):
return self._flattened_module.training and torch.is_grad_enabled()
return self.training and torch.is_grad_enabled()
def train(self: T, mode: bool = True) -> T:
"""Override original method to delegate execution to the flattened PyTorch user module"""
# Since _modules is empty, the task needs to be delegated to _module.flattened_module.train
# which will recursively update the original_module
self.training = mode
self._module_metadata.flattened_module.train(mode)
return self
def state_dict(self, destination=None, prefix='', keep_vars=False):
"""Override original method to delegate execution to the base module"""
"""Override original method to delegate execution to the original PyTorch user module"""
# Override the state_dict() method so that the state dict key names
# do not contain the _flattened_module._original_module prefix
return self._original_module.state_dict(
# do not contain the flattened_module._original_module prefix
return self._module_metadata.original_module.state_dict(
destination=destination, prefix=prefix, keep_vars=keep_vars)
def load_state_dict(self, state_dict: 'OrderedDict[str, Tensor]',
strict: bool = True):
"""Override original method to delegate execution to the base module"""
"""Override original method to delegate execution to the original PyTorch user module"""
# Override the load_state_dict() method so that the loaded state dict
# key names does not need to contain the _flattened_module._original_module prefix
return self._original_module.load_state_dict(
# key names does not need to contain the _module.flattened_module._original_module prefix
return self._module_metadata.original_module.load_state_dict(
state_dict, strict=strict)
def register_buffer(self, name: str, tensor: Optional[torch.Tensor], persistent: bool = True) -> None:
"""Override original method to delegate execution to the base module"""
self._original_module.register_buffer(name, tensor, persistent=persistent)
"""Override original method to delegate execution to the original PyTorch user module"""
self._module_metadata.original_module.register_buffer(name, tensor, persistent=persistent)
def register_parameter(self, name: str, param: Optional[torch.nn.Parameter]) -> None:
"""Override original method to delegate execution to the base module"""
self._original_module.register_parameter(name, param)
"""Override original method to delegate execution to the original PyTorch user module"""
self._module_metadata.original_module.register_parameter(name, param)
def get_parameter(self, target: str) -> torch.nn.Parameter:
"""Override original method to delegate execution to the base module"""
return self._original_module.get_parameter(target)
"""Override original method to delegate execution to the original PyTorch user module"""
return self._module_metadata.original_module.get_parameter(target)
def get_buffer(self, target: str) -> torch.Tensor:
"""Override original method to delegate execution to the base module"""
return self._original_module.get_buffer(target)
"""Override original method to delegate execution to the original PyTorch user module"""
return self._module_metadata.original_module.get_buffer(target)
def parameters(self, recurse: bool = True) -> Iterator[torch.nn.Parameter]:
"""Override original method to delegate execution to the base module"""
yield from self._original_module.parameters(recurse=recurse)
"""Override original method to delegate execution to the original PyTorch user module"""
yield from self._module_metadata.original_module.parameters(recurse=recurse)
def named_parameters(self, prefix: str = '', recurse: bool = True) -> Iterator[Tuple[str, torch.nn.Parameter]]:
"""Override original method to delegate execution to the base module"""
yield from self._original_module.named_parameters(prefix=prefix, recurse=recurse)
"""Override original method to delegate execution to the original PyTorch user module"""
yield from self._module_metadata.original_module.named_parameters(prefix=prefix, recurse=recurse)
def buffers(self, recurse: bool = True) -> Iterator[torch.Tensor]:
"""Override original method to delegate execution to the base module"""
yield from self._original_module.buffers(recurse=recurse)
"""Override original method to delegate execution to the original PyTorch user module"""
yield from self._module_metadata.original_module.buffers(recurse=recurse)
def named_buffers(self, prefix: str = '', recurse: bool = True) -> Iterator[Tuple[str, torch.Tensor]]:
"""Override original method to delegate execution to the base module"""
yield from self._original_module.named_buffers(prefix=prefix, recurse=recurse)
"""Override original method to delegate execution to the original PyTorch user module"""
yield from self._module_metadata.original_module.named_buffers(prefix=prefix, recurse=recurse)
def _replicate_for_data_parallel(self):
"""Raises a NotImplementedError exception since ORTModule is not compatible with torch.nn.DataParallel
@ -135,3 +160,34 @@ class ORTModule(torch.nn.Module):
raise NotImplementedError("ORTModule is not compatible with torch.nn.DataParallel. "
"Please use torch.nn.parallel.DistributedDataParallel instead.")
def _load_from_state_dict(self, state_dict, prefix, local_metadata, strict,
missing_keys, unexpected_keys, error_msgs):
"""Override original method to delegate execution to the original PyTorch user module"""
# PyTorch load_state_dict implementation does not recursively call load_state_dict on its sub-modules.
# Instead, it creates a recursive function and invokes _load_from_state_dict on all child modules.
# For the scenario where an ORTModule is a sub-module of another module, loading of the state
# dictionary requires the _load_from_state_dict to be overridden to prevent an error.
self._module_metadata.original_module._load_from_state_dict(state_dict, prefix, local_metadata, strict,
missing_keys, unexpected_keys, error_msgs)
def named_children(self) -> Iterator[Tuple[str, 'Module']]:
"""Override original method to delegate execution to the original PyTorch user module"""
yield from self._module_metadata.original_module.named_children()
def modules(self) -> Iterator['Module']:
"""Override original method to delegate execution to the original PyTorch user module"""
yield from self._module_metadata.original_module.modules()
def named_modules(self, memo: Optional[Set['Module']] = None, prefix: str = ''):
"""Override original method to delegate execution to the original PyTorch user module"""
yield from self._module_metadata.original_module.named_modules(memo, prefix)
def add_module(self, name: str, module: Optional['Module']) -> None:
"""Raises a NotImplementedError exception since ORTModule does not support adding modules to it"""
raise NotImplementedError("ORTModule does not support adding modules to it.")

View file

@ -592,7 +592,7 @@ def test_gradient_correctness_conv1d(use_fp16, input_requires_grad):
if use_fp16:
_test_helpers.assert_values_are_close(ort_prediction, pt_prediction, atol=1e-3, rtol=1e-3)
_test_helpers.assert_gradients_match_and_reset_gradient(ort_model, pt_model, rtol=1e-2, atol=1.1e-2)
_test_helpers.assert_gradients_match_and_reset_gradient(ort_model, pt_model, rtol=1e-2, atol=2e-2)
else:
_test_helpers.assert_values_are_close(ort_prediction, pt_prediction, atol=1e-5)
_test_helpers.assert_gradients_match_and_reset_gradient(ort_model, pt_model, rtol=5e-3, atol=4e-3)
@ -1666,26 +1666,26 @@ def test_model_initializer_requires_grad_changes_from_one_forward_to_next():
model.fc1.requires_grad_(True)
model = ORTModule(model)
x = torch.randn(N, D_in, device=device)
assert model._original_module.fc1.weight.grad is None
assert model._original_module.fc1.bias.grad is None
assert model._module_metadata.original_module.fc1.weight.grad is None
assert model._module_metadata.original_module.fc1.bias.grad is None
# Make sure no exception is raised
output = model(x)
loss = torch.sum(output)
loss.backward()
training_session1 = model._execution_manager(model._is_training())._execution_agent
weight_grad_2 = model._original_module.fc1.weight.grad
bias_grad_2 = model._original_module.fc1.bias.grad
weight_grad_2 = model._module_metadata.original_module.fc1.weight.grad
bias_grad_2 = model._module_metadata.original_module.fc1.bias.grad
assert weight_grad_2 is not None
assert bias_grad_2 is not None
model._original_module.fc1.requires_grad_(False)
model._module_metadata.original_module.fc1.requires_grad_(False)
output = model(x)
loss = torch.sum(output)
loss.backward()
training_session2 = model._execution_manager(model._is_training())._execution_agent
weight_grad_3 = model._original_module.fc1.weight.grad
bias_grad_3 = model._original_module.fc1.bias.grad
weight_grad_3 = model._module_metadata.original_module.fc1.weight.grad
bias_grad_3 = model._module_metadata.original_module.fc1.bias.grad
assert training_session1 != training_session2
assert torch.equal(weight_grad_2, weight_grad_3)
@ -2619,3 +2619,31 @@ def test_unused_parameters_does_not_unnecssarily_reinitilize(model):
{})
assert not training_manager._reinitialize_graph_builder(input_info)
def test_load_state_dict_for_wrapped_ortmodule():
class WrapperModule(torch.nn.Module):
def __init__(self, ortmodule):
super(WrapperModule, self).__init__()
self._ortmodule = ortmodule
def forward(self, x):
return self._ortmodule(x)
device = 'cuda'
N, D_in, H, D_out = 64, 784, 500, 10
model = NeuralNetSinglePositionalArgument(D_in, H, D_out).to(device)
model = ORTModule(copy.deepcopy(model))
wrapper_module = WrapperModule(model)
x = torch.randn(N, D_in, device=device)
_ = wrapper_module(x)
state_dict1 = wrapper_module.state_dict()
list(next(iter(state_dict1.items())))[1] += 10
wrapper_module.load_state_dict(state_dict1)
state_dict2 = wrapper_module.state_dict()
assert state_dict1
assert len(state_dict1.keys()) == len(state_dict2.keys())
for param_name, param_value in state_dict1.items():
assert param_name in state_dict2
assert torch.equal(param_value, state_dict2[param_name])

View file

@ -266,6 +266,8 @@ requirements_file = "requirements.txt"
local_version = None
enable_training = parse_arg_remove_boolean(sys.argv, '--enable_training')
default_training_package_device = parse_arg_remove_boolean(sys.argv, '--default_training_package_device')
if enable_training:
packages.extend(['onnxruntime.training',
'onnxruntime.training.amp',
@ -280,13 +282,16 @@ if enable_training:
# this is needed immediately by pytorch/ort so that the user is able to
# install an onnxruntime training package with matching torch cuda version.
package_name = 'onnxruntime-training'
if cuda_version:
# removing '.' to make local Cuda version number in the same form as Pytorch.
local_version = '+cu' + cuda_version.replace('.', '')
if rocm_version:
# removing '.' to make Cuda version number in the same form as Pytorch.
rocm_version = rocm_version.replace('.', '')
local_version = '+rocm' + rocm_version
# we want put default training packages to pypi. pypi does not accept package with a local version.
if not default_training_package_device:
if cuda_version:
# removing '.' to make local Cuda version number in the same form as Pytorch.
local_version = '+cu' + cuda_version.replace('.', '')
if rocm_version:
# removing '.' to make Cuda version number in the same form as Pytorch.
rocm_version = rocm_version.replace('.', '')
local_version = '+rocm' + rocm_version
package_data = {}

View file

@ -1536,7 +1536,8 @@ def run_nodejs_tests(nodejs_binding_dir):
def build_python_wheel(
source_dir, build_dir, configs, use_cuda, cuda_version, use_rocm, rocm_version, use_dnnl,
use_tensorrt, use_openvino, use_nuphar, use_vitisai, use_acl, use_armnn, use_dml,
wheel_name_suffix, enable_training, nightly_build=False, featurizers_build=False, use_ninja=False):
wheel_name_suffix, enable_training, nightly_build=False, default_training_package_device=False,
featurizers_build=False, use_ninja=False):
for config in configs:
cwd = get_config_build_dir(build_dir, config)
if is_windows() and not use_ninja:
@ -1558,6 +1559,8 @@ def build_python_wheel(
# Any combination of the following arguments can be applied
if nightly_build:
args.append('--nightly_build')
if default_training_package_device:
args.append('--default_training_package_device')
if featurizers_build:
args.append("--use_featurizers")
if wheel_name_suffix:
@ -2087,6 +2090,7 @@ def main():
if args.build:
if args.build_wheel:
nightly_build = bool(os.getenv('NIGHTLY_BUILD') == '1')
default_training_package_device = bool(os.getenv('DEFAULT_TRAINING_PACKAGE_DEVICE') == '1')
build_python_wheel(
source_dir,
build_dir,
@ -2106,6 +2110,7 @@ def main():
args.wheel_name_suffix,
args.enable_training,
nightly_build=nightly_build,
default_training_package_device=default_training_package_device,
featurizers_build=args.use_featurizers,
use_ninja=(args.cmake_generator == 'Ninja')
)

View file

@ -72,6 +72,29 @@ def _parse_build_settings(args):
return build_settings
# Add ORT C and C++ API headers to the AAR package (in fact a zip file)
# Such that developers using ORT native API can extract libraries and header from AAR package without building ORT
# TODO, see if we can use Gradle to add headers to AAR package directly, which is necessary if we want to
# publish the packagee directly using Gradle in the pipeline
def _add_headers_to_aar(aar_file_path, header_files_path):
import shutil
import tempfile
with tempfile.TemporaryDirectory() as temp_dir:
aar_content = os.path.join(temp_dir, 'aar_content')
shutil.unpack_archive(aar_file_path, aar_content, 'zip')
# copy necessary header files
shutil.copytree(header_files_path, os.path.join(aar_content, 'headers'))
# create the zip archive
zip_base_filename = os.path.join(temp_dir, 'aar_with_headers')
zip_filename = zip_base_filename + '.zip'
shutil.make_archive(zip_base_filename, 'zip', root_dir=aar_content)
# overwrite the existing AAR package
shutil.move(zip_filename, aar_file_path)
def _build_aar(args):
build_settings = _parse_build_settings(args)
build_dir = os.path.abspath(args.build_dir)
@ -89,6 +112,7 @@ def _build_aar(args):
_base_build_command = [
sys.executable, BUILD_PY, '--config=' + _build_config
] + build_settings['build_params']
header_files_path = ''
# Build binary for each ABI, one by one
for abi in build_settings['build_abis']:
@ -116,6 +140,10 @@ def _build_aar(args):
os.remove(_target_lib_name)
os.symlink(os.path.join(_build_dir, _build_config, lib_name), _target_lib_name)
# we only need to define the header files path once
if not header_files_path:
header_files_path = os.path.join(_build_dir, _build_config, 'android', 'headers')
# The directory to publish final AAR
_aar_publish_dir = os.path.join(build_dir, 'aar_out', _build_config)
os.makedirs(_aar_publish_dir, exist_ok=True)
@ -139,6 +167,11 @@ def _build_aar(args):
# clean, build, and publish to a local directory
subprocess.run(_gradle_command + ['clean'], env=_env, shell=_shell, check=True, cwd=JAVA_ROOT)
subprocess.run(_gradle_command + ['build'], env=_env, shell=_shell, check=True, cwd=JAVA_ROOT)
# add C and C++ API headers to the intermediate aar package
aar_file_path = os.path.join(_aar_dir, 'outputs', 'aar', 'onnxruntime-release.aar')
_add_headers_to_aar(aar_file_path, header_files_path)
subprocess.run(_gradle_command + ['publish'], env=_env, shell=_shell, check=True, cwd=JAVA_ROOT)

View file

@ -119,7 +119,7 @@ stages:
- task: CmdLine@2
displayName: 'Build Python Documentation'
condition: ne(variables['PythonVersion'], '3.9') # tensorflow not available on python 3.9
condition: and(succeeded(), ne(variables['PythonVersion'], '3.9')) # tensorflow not available on python 3.9
inputs:
script: |
mkdir -p $HOME/.onnx
@ -137,7 +137,7 @@ stages:
- task: CopyFiles@2
displayName: 'Copy Python Documentation to: $(Build.ArtifactStagingDirectory)'
condition: ne(variables['PythonVersion'], '3.9') # tensorflow not available on python 3.9
condition: and(succeeded(), ne(variables['PythonVersion'], '3.9')) # tensorflow not available on python 3.9
inputs:
SourceFolder: '$(Build.BinariesDirectory)/docs/inference/html'
Contents: '**'
@ -431,7 +431,7 @@ stages:
- task: CmdLine@2
displayName: 'Build Python Documentation'
condition: ne(variables['PythonVersion'], '3.9') # tensorflow not available on python 3.9
condition: and(succeeded(), ne(variables['PythonVersion'], '3.9')) # tensorflow not available on python 3.9
inputs:
script: |
mkdir -p $HOME/.onnx
@ -447,7 +447,7 @@ stages:
- task: CopyFiles@2
displayName: 'Copy Python Documentation to: $(Build.ArtifactStagingDirectory)'
condition: ne(variables['PythonVersion'], '3.9') # tensorflow not available on python 3.9
condition: and(succeeded(), ne(variables['PythonVersion'], '3.9')) # tensorflow not available on python 3.9
inputs:
SourceFolder: '$(Build.BinariesDirectory)/docs/training/html'
Contents: '**'
@ -545,6 +545,7 @@ stages:
--volume $HOME/.onnx:/home/onnxruntimedev/.onnx \
-e NVIDIA_VISIBLE_DEVICES=all \
-e NIGHTLY_BUILD \
-e DEFAULT_TRAINING_PACKAGE_DEVICE \
-e BUILD_BUILDNUMBER \
onnxruntimetraininggpubuild \
$(PythonManylinuxDir)/bin/python3 /onnxruntime_src/tools/ci_build/build.py \
@ -588,7 +589,7 @@ stages:
- task: CmdLine@2
displayName: 'Build Python Documentation'
condition: ne(variables['PythonVersion'], '3.9') # tensorflow not available on python 3.9
condition: and(succeeded(), ne(variables['PythonVersion'], '3.9')) # tensorflow not available on python 3.9
inputs:
script: |
mkdir -p $HOME/.onnx
@ -606,7 +607,7 @@ stages:
- task: CopyFiles@2
displayName: 'Copy Python Documentation to: $(Build.ArtifactStagingDirectory)'
condition: ne(variables['PythonVersion'], '3.9') # tensorflow not available on python 3.9
condition: and(succeeded(), ne(variables['PythonVersion'], '3.9')) # tensorflow not available on python 3.9
inputs:
SourceFolder: '$(Build.BinariesDirectory)/docs/training/html'
Contents: '**'
@ -761,7 +762,7 @@ stages:
- task: CmdLine@2
displayName: 'Build Python Documentation'
condition: ne(variables['PythonVersion'], '3.9') # tensorflow not available on python 3.9
condition: and(succeeded(), ne(variables['PythonVersion'], '3.9')) # tensorflow not available on python 3.9
inputs:
script: |
mkdir -p $HOME/.onnx
@ -779,7 +780,7 @@ stages:
- task: CopyFiles@2
displayName: 'Copy Python Documentation to: $(Build.ArtifactStagingDirectory)'
condition: ne(variables['PythonVersion'], '3.9') # tensorflow not available on python 3.9
condition: and(succeeded(), ne(variables['PythonVersion'], '3.9')) # tensorflow not available on python 3.9
inputs:
SourceFolder: '$(Build.BinariesDirectory)/docs/training/html'
Contents: '**'

View file

@ -189,9 +189,17 @@ jobs:
workingDirectory: '$(Build.SourcesDirectory)\js'
displayName: 'Clang-format'
- script: |
node -e "a=require('child_process').execSync('git ls-files -m').toString();if(a)throw new Error('Following source files are not formatted:\n'+a)"
node -e "a=require('child_process').execSync('git ls-files -m').toString();if(a)throw new Error('Following source files are not formatted: (did you run \"npm run format\"?)\n'+a)"
workingDirectory: '$(Build.SourcesDirectory)\js'
displayName: 'Check unformatted files'
- script: |
npm run build:doc
workingDirectory: '$(Build.SourcesDirectory)\js\web'
displayName: 'Generating documents'
- script: |
node -e "a=require('child_process').execSync('git ls-files -m').toString();if(a)throw new Error('Following documents are not up-to-date: (did you run \"npm run build:doc\"?)\n'+a)"
workingDirectory: '$(Build.SourcesDirectory)\js\web'
displayName: 'Check out of dated documents'
- script: |
npm run build
workingDirectory: '$(Build.SourcesDirectory)\js\web'
@ -199,7 +207,11 @@ jobs:
- script: |
npm test
workingDirectory: '$(Build.SourcesDirectory)\js\web'
displayName: 'Run ort-web tests'
displayName: 'Run ort-web tests - unpacked mode'
- script: |
npm test -- --webgl-texture-pack-mode -b=webgl
workingDirectory: '$(Build.SourcesDirectory)\js\web'
displayName: 'Run ort-web tests - packed mode'
- script: |
npm pack
workingDirectory: '$(Build.SourcesDirectory)\js\common'