mirror of
https://github.com/saymrwulf/onnxruntime.git
synced 2026-07-07 04:39:07 +00:00
Migraphx improvements (#4328)
* Add amd migraphx execution provider to onnx runtime * rename MiGraphX to MIGraphX * add migraphx EP to tests * support multiple program output * disable more tests * backup changes related to program multiple outputs * remove logging code * remove unnecessary changes in migraphx_execution_provider.cc * add migraphx EP to tests * add input requests of the batchnorm operator * add to support an onnx operator PRelu * update migrapx dockerfile and removed one unused line * chagnes related to support dynamic input shape * fix build error * code backup * code backup * version that has 106 models run correctly * code backup * code backup * remove unnecessary print info * code backup * code backup * code backup * code backup * code backup * code backup * changes corresponding to migraphx change * fix merge conflict * minor code cleanup * code cleanup * remove unnecessary code * remove unnecessary code * add to support more constant folding analysis * more constant folding checking for shape input * add env var to control whether fp16 is enabled. Modify docker file to use ROCM3.3 * fix function name to avoid build error * add build and execution instruction for migraphx execution provider * added more build instructions * fixed a small format error * a minor change * fix review comments * another minor change * additional refinement of the documents * additional changes * remove unnecessary changes in the dockfile * additional changes for the dockerfile * code change backup * fix errors related to a few unit tests * fix a build error related to api change * fix unit test errors by either disabling the test or fix related isssues * remove unnecessary log info * sync submodule tvm with master * remove unnecessary changes * remove an unnecessary code line * refine documents for addition example
This commit is contained in:
Shucai Xiao
GitHub
parent
0b450dcd9f
commit
bfc888613f
11 changed files with 270 additions and 144 deletions
24
BUILD.md
24
BUILD.md
|
|
@ -22,6 +22,7 @@
|
|||
* [ArmNN](#ArmNN)
|
||||
* [Rockchip RKNPU](#RKNPU)
|
||||
* [Xilinx Vitis-AI](#Vitis-AI)
|
||||
* [AMD MIGraphX](#AMD-MIGraphX)
|
||||
* Options
|
||||
* [OpenMP](#OpenMP)
|
||||
* [OpenBLAS](#OpenBLAS)
|
||||
|
|
@ -207,6 +208,7 @@ See more information on the TensorRT Execution Provider [here](./docs/execution_
|
|||
|
||||
Dockerfile instructions are available [here](./dockerfiles#tensorrt)
|
||||
|
||||
---
|
||||
|
||||
#### Jetson TX1/TX2/Nano (ARM64 Builds)
|
||||
|
||||
|
|
@ -989,6 +991,28 @@ Android Archive (AAR) files, which can be imported directly in Android Studio, w
|
|||
|
||||
If you want to use NNAPI Execution Provider on Android, see [docs/execution_providers/NNAPI-ExecutionProvider.md](/docs/execution_providers/NNAPI-ExecutionProvider.md).
|
||||
|
||||
---
|
||||
|
||||
### AMD MIGraphX
|
||||
|
||||
See more information on the MIGraphX Execution Provider [here](./docs/execution_providers/MIGraphX-ExecutionProvider.md).
|
||||
|
||||
#### Prerequisites
|
||||
* Install [ROCM](https://rocmdocs.amd.com/en/latest/Installation_Guide/Installation-Guide.html)
|
||||
* The MIGraphX execution provider for ONNX Runtime is built and tested with ROCM3.3
|
||||
* Install [MIGraphX](https://github.com/ROCmSoftwarePlatform/AMDMIGraphX)
|
||||
* The path to MIGraphX installation must be provided via the `--migraphx_home parameter`.
|
||||
|
||||
#### Build Instructions
|
||||
|
||||
##### Linux
|
||||
|
||||
```
|
||||
./build.sh --config <Release|Debug|RelWithDebInfo> --use_migraphx --migraphx_home <path to MIGraphX home>
|
||||
```
|
||||
|
||||
Dockerfile instructions are available [here](./dockerfiles#migraphx)
|
||||
|
||||
***
|
||||
|
||||
# Training
|
||||
|
|
|
|||
|
|
@ -137,7 +137,7 @@ For production scenarios, it's strongly recommended to build only from an [offic
|
|||
|
||||
|CPU|GPU|IoT/Edge/Mobile|Other|
|
||||
|---|---|---|---|
|
||||
|<ul><li>Default CPU - *MLAS (Microsoft Linear Algebra Subprograms) + Eigen*</li><li>[Intel DNNL](./docs/execution_providers/DNNL-ExecutionProvider.md)</li><li>[Intel nGraph](./docs/execution_providers/nGraph-ExecutionProvider.md)</li><li>Intel MKL-ML *(build option)*</li></ul>|<ul><li>NVIDIA CUDA</li><li>[NVIDIA TensorRT](./docs/execution_providers/TensorRT-ExecutionProvider.md)</li><li>[DirectML](./docs/execution_providers/DirectML-ExecutionProvider.md)</li></ul>|<ul><li>[Intel OpenVINO](./docs/execution_providers/OpenVINO-ExecutionProvider.md)</li><li>[ARM Compute Library](./docs/execution_providers/ACL-ExecutionProvider.md) (*preview*)</li><li>[Android Neural Networks API](./docs/execution_providers/NNAPI-ExecutionProvider.md) (*preview*)</li></ul>|<ul><li>[Nuphar Model Compiler](./docs/execution_providers/Nuphar-ExecutionProvider.md) - (*preview*)</li><li>[Rockchip NPU](./docs/execution_providers/RKNPU-ExecutionProvider.md) (*preview*)</li><li>[Xilinx Vitis-AI](./docs/execution_providers/Vitis-AI-ExecutionProvider.md) (*preview*)</li></ul>|
|
||||
|<ul><li>Default CPU - *MLAS (Microsoft Linear Algebra Subprograms) + Eigen*</li><li>[Intel DNNL](./docs/execution_providers/DNNL-ExecutionProvider.md)</li><li>[Intel nGraph](./docs/execution_providers/nGraph-ExecutionProvider.md)</li><li>Intel MKL-ML *(build option)*</li></ul>|<ul><li>NVIDIA CUDA</li><li>[NVIDIA TensorRT](./docs/execution_providers/TensorRT-ExecutionProvider.md)</li><li>[DirectML](./docs/execution_providers/DirectML-ExecutionProvider.md)</li><li>[AMD MIGraphX](./docs/execution_providers/MIGraphX-ExecutionProvider.md)</li></ul>|<ul><li>[Intel OpenVINO](./docs/execution_providers/OpenVINO-ExecutionProvider.md)</li><li>[ARM Compute Library](./docs/execution_providers/ACL-ExecutionProvider.md) (*preview*)</li><li>[Android Neural Networks API](./docs/execution_providers/NNAPI-ExecutionProvider.md) (*preview*)</li></ul>|<ul><li>[Nuphar Model Compiler](./docs/execution_providers/Nuphar-ExecutionProvider.md) - (*preview*)</li><li>[Rockchip NPU](./docs/execution_providers/RKNPU-ExecutionProvider.md) (*preview*)</li><li>[Xilinx Vitis-AI](./docs/execution_providers/Vitis-AI-ExecutionProvider.md) (*preview*)</li></ul>|
|
||||
|
||||
* [Roadmap: Upcoming accelerators](./docs/Roadmap.md#accelerators-and-execution-providers)
|
||||
* [Extensibility: Add an execution provider](docs/AddingExecutionProvider.md)
|
||||
|
|
|
|||
|
|
@ -10,14 +10,18 @@ FROM ubuntu:16.04
|
|||
ARG ONNXRUNTIME_REPO=https://github.com/Microsoft/onnxruntime
|
||||
ARG ONNXRUNTIME_BRANCH=master
|
||||
ENV DEBIAN_FRONTEND noninteractive
|
||||
ENV MIGRAPHX_DISABLE_FAST_GELU=1
|
||||
|
||||
RUN apt-get clean && apt-get update && apt-get install -y locales
|
||||
RUN locale-gen en_US.UTF-8
|
||||
RUN update-locale LANG=en_US.UTF-8
|
||||
ENV LC_ALL C.UTF-8
|
||||
ENV LANG C.UTF-8
|
||||
ENV MIGRAPHX_DISABLE_FAST_GELU=1
|
||||
|
||||
# Install rocm
|
||||
RUN apt-get update && apt-get install -y --no-install-recommends curl && \
|
||||
curl -sL http://repo.radeon.com/rocm/apt/debian/rocm.gpg.key | apt-key add - && \
|
||||
sh -c 'echo deb [arch=amd64] http://repo.radeon.com/rocm/apt/debian/ xenial main > /etc/apt/sources.list.d/rocm.list'
|
||||
sh -c 'echo deb [arch=amd64] http://repo.radeon.com/rocm/apt/3.3/ xenial main > /etc/apt/sources.list.d/rocm.list'
|
||||
|
||||
RUN apt-get update &&\
|
||||
apt-get install -y sudo git bash build-essential cmake libpython3.5-dev python3-pip miopen-hip rocblas half
|
||||
|
|
|
|||
|
|
@ -12,6 +12,7 @@
|
|||
- ARM 32v7: [Dockerfile](Dockerfile.arm32v7), [Instructions](#arm-32v7)
|
||||
- ONNX-Ecosystem (CPU + Converters): [Dockerfile](https://github.com/onnx/onnx-docker/blob/master/onnx-ecosystem/Dockerfile), [Instructions](https://github.com/onnx/onnx-docker/tree/master/onnx-ecosystem)
|
||||
- ONNX Runtime Server: [Dockerfile](Dockerfile.server), [Instructions](#onnx-runtime-server)
|
||||
- MIGraphX: [Dockerfile](Dockerfile.migraphx), [Instructions](#migraphx)
|
||||
|
||||
**Published Microsoft Container Registry (MCR) Images**
|
||||
|
||||
|
|
@ -27,10 +28,11 @@ Use `docker pull` with any of the images and tags below to pull an image and try
|
|||
| OpenVino (VAD-M) | mcr.microsoft.com/azureml/onnxruntime | :v0.5.0-openvino-r1.1-vadm, :v1.0.0-openvino-r1.1-vadm | :latest-openvino-vadm |
|
||||
| OpenVino (MYRIAD) | mcr.microsoft.com/azureml/onnxruntime | :v0.5.0-openvino-r1.1-myriad, :v1.0.0-openvino-r1.1-myriad, :v1.3.0-openvino-2020.2.120-myriad| :latest-openvino-myriad |
|
||||
| OpenVino (CPU) | mcr.microsoft.com/azureml/onnxruntime | :v1.0.0-openvino-r1.1-cpu, :v1.3.0-openvino-2020.2.120-cpu | :latest-openvino-cpu |
|
||||
| OpenVINO (GPU) | mcr.microsoft.com/azureml/onnxruntime | :v1.3.0-openvino-2020.2.120-gpu | :latest-openvino-gpu |
|
||||
| OpenVINO (GPU) | mcr.microsoft.com/azureml/onnxruntime | :v1.3.0-openvino-2020.2.120-gpu | :latest-openvino-gpu |
|
||||
| nGraph | mcr.microsoft.com/azureml/onnxruntime | :v1.0.0-ngraph-v0.26.0 | :latest-ngraph |
|
||||
| Nuphar | mcr.microsoft.com/azureml/onnxruntime | | :latest-nuphar |
|
||||
| Server | mcr.microsoft.com/onnxruntime/server | :v0.4.0, :v0.5.0, :v0.5.1, :v1.0.0 | :latest |
|
||||
| MIGraphX (GPU) | mcr.microsoft.com/azureml/onnxruntime | :v0.6 | :latest |
|
||||
| Training ([usage](https://github.com/microsoft/onnxruntime-training-examples))| mcr.microsoft.com/azureml/onnxruntime-training | :0.1-rc1-openmpi4.0-cuda10.1-cudnn7.6-nccl2.4.8| 0.1-rc1-openmpi4.0-cuda10.1-cudnn7.6-nccl2.4.8|
|
||||
---
|
||||
|
||||
|
|
@ -248,6 +250,20 @@ The Dockerfile used in these instructions specifically targets Raspberry Pi 3/3+
|
|||
docker run -it onnxruntime-nuphar
|
||||
```
|
||||
|
||||
## MIGraphX
|
||||
**Ubuntu 16.04, rocm3.3, AMDMIGraphX v0.7**
|
||||
|
||||
1. Build the docker image from the Dockerfile in this repository.
|
||||
```
|
||||
docker build -t onnxruntime-migraphx -f Dockerfile.migraphx .
|
||||
```
|
||||
|
||||
2. Run the Docker image
|
||||
|
||||
```
|
||||
docker run -it --device=/dev/kfd --device=/dev/dri --group-add video onnxruntime-migraphx
|
||||
```
|
||||
|
||||
## ONNX Runtime Server
|
||||
*Public Preview*
|
||||
|
||||
|
|
|
|||
|
|
@ -61,7 +61,6 @@ To achieve the best performance on a growing set of compute targets across cloud
|
|||
Supported
|
||||
|
||||
Supported EPs are listed [here](../README.md#supported-accelerators). Upcoming EPs include:
|
||||
* AMD GPU
|
||||
* Xilinx FPGA
|
||||
|
||||
|
||||
|
|
|
|||
43
docs/execution_providers/MIGraphX-ExecutionProvider.md
Normal file
43
docs/execution_providers/MIGraphX-ExecutionProvider.md
Normal file
|
|
@ -0,0 +1,43 @@
|
|||
# MIGraphX Execution Provider
|
||||
|
||||
ONNX Runtime's [MIGraphX](https://github.com/ROCmSoftwarePlatform/AMDMIGraphX/) execution provider uses AMD's Deep Learning graph optimization engine to accelerate ONNX model on AMD GPUs.
|
||||
|
||||
## Build
|
||||
For build instructions, please see the [BUILD page](../../BUILD.md#AMD-MIGraphX).
|
||||
|
||||
## Using the MIGraphX execution provider
|
||||
### C/C++
|
||||
The MIGraphX execution provider needs to be registered with ONNX Runtime to enable in the inference session.
|
||||
```
|
||||
string log_id = "Foo";
|
||||
auto logging_manager = std::make_unique<LoggingManager>
|
||||
(std::unique_ptr<ISink>{new CLogSink{}},
|
||||
static_cast<Severity>(lm_info.default_warning_level),
|
||||
false,
|
||||
LoggingManager::InstanceType::Default,
|
||||
&log_id)
|
||||
Environment::Create(std::move(logging_manager), env)
|
||||
InferenceSession session_object{so,env};
|
||||
session_object.RegisterExecutionProvider(std::make_unique<::onnxruntime::MIGraphXExecutionProvider>());
|
||||
status = session_object.Load(model_file_name);
|
||||
```
|
||||
You can check [here](https://github.com/scxiao/ort_test/tree/master/char_rnn) for a specific c/c++ program.
|
||||
|
||||
The C API details are [here](../C_API.md#c-api).
|
||||
|
||||
### Python
|
||||
When using the Python wheel from the ONNX Runtime build with MIGraphX execution provider, it will be automatically
|
||||
prioritized over the default GPU or CPU execution providers. There is no need to separately register the execution
|
||||
provider. Python APIs details are [here](../python/api_summary.rst#api-summary).
|
||||
|
||||
You can check [here](https://github.com/scxiao/ort_test/tree/master/python/run_onnx) for a python script to run an
|
||||
model on either the CPU or MIGraphX Execution Provider.
|
||||
|
||||
## Performance Tuning
|
||||
For performance tuning, please see guidance on this page: [ONNX Runtime Perf Tuning](../ONNX_Runtime_Perf_Tuning.md)
|
||||
|
||||
When/if using [onnxruntime_perf_test](../../onnxruntime/test/perftest#onnxruntime-performance-test), use the flag `-e migraphx`
|
||||
|
||||
## Configuring environment variables
|
||||
MIGraphX providers an environment variable ORT_MIGRAPHX_FP16_ENABLE to enable the FP16 mode.
|
||||
|
||||
|
|
@ -40,10 +40,6 @@ void HIPAllocator::Free(void* p) {
|
|||
hipFree(p); // do not throw error since it's OK for hipFree to fail during shutdown
|
||||
}
|
||||
|
||||
const OrtMemoryInfo& HIPAllocator::Info() const {
|
||||
return info_;
|
||||
}
|
||||
|
||||
FencePtr HIPAllocator::CreateFence(const SessionState* session_state) {
|
||||
return std::make_shared<HIPFence>(GetGPUDataTransfer(session_state));
|
||||
}
|
||||
|
|
@ -60,10 +56,6 @@ void HIPPinnedAllocator::Free(void* p) {
|
|||
hipHostFree(p);
|
||||
}
|
||||
|
||||
const OrtMemoryInfo& HIPPinnedAllocator::Info() const {
|
||||
return info_;
|
||||
}
|
||||
|
||||
FencePtr HIPPinnedAllocator::CreateFence(const SessionState* session_state) {
|
||||
return std::make_shared<HIPFence>(GetGPUDataTransfer(session_state));
|
||||
}
|
||||
|
|
|
|||
|
|
@ -9,30 +9,32 @@ namespace onnxruntime {
|
|||
|
||||
class HIPAllocator : public IDeviceAllocator {
|
||||
public:
|
||||
HIPAllocator(int device_id, const char* name) : info_(name, OrtAllocatorType::OrtDeviceAllocator, OrtDevice(OrtDevice::GPU, OrtDevice::MemType::DEFAULT, device_id), device_id, OrtMemTypeDefault) {}
|
||||
HIPAllocator(int device_id, const char* name)
|
||||
: IDeviceAllocator(
|
||||
OrtMemoryInfo(name, OrtAllocatorType::OrtDeviceAllocator,
|
||||
OrtDevice(OrtDevice::GPU, OrtDevice::MemType::DEFAULT, device_id),
|
||||
device_id, OrtMemTypeDefault)) {}
|
||||
|
||||
virtual void* Alloc(size_t size) override;
|
||||
virtual void Free(void* p) override;
|
||||
virtual const OrtMemoryInfo& Info() const override;
|
||||
virtual FencePtr CreateFence(const SessionState* session_state) override;
|
||||
|
||||
private:
|
||||
void CheckDevice() const;
|
||||
|
||||
private:
|
||||
const OrtMemoryInfo info_;
|
||||
};
|
||||
|
||||
//TODO: add a default constructor
|
||||
class HIPPinnedAllocator : public IDeviceAllocator {
|
||||
public:
|
||||
HIPPinnedAllocator(int device_id, const char* name) : info_(name, OrtAllocatorType::OrtDeviceAllocator, OrtDevice(OrtDevice::CPU, OrtDevice::MemType::HIP_PINNED, device_id), device_id, OrtMemTypeCPUOutput) {}
|
||||
HIPPinnedAllocator(int device_id, const char* name)
|
||||
: IDeviceAllocator(
|
||||
OrtMemoryInfo(name, OrtAllocatorType::OrtDeviceAllocator,
|
||||
OrtDevice(OrtDevice::CPU, OrtDevice::MemType::HIP_PINNED, device_id),
|
||||
device_id, OrtMemTypeCPUOutput)) {}
|
||||
|
||||
virtual void* Alloc(size_t size) override;
|
||||
virtual void Free(void* p) override;
|
||||
virtual const OrtMemoryInfo& Info() const override;
|
||||
virtual FencePtr CreateFence(const SessionState* session_state) override;
|
||||
|
||||
private:
|
||||
const OrtMemoryInfo info_;
|
||||
};
|
||||
|
||||
} // namespace onnxruntime
|
||||
|
|
|
|||
|
|
@ -10,6 +10,7 @@
|
|||
#include "core/graph/graph_viewer.h"
|
||||
#include "core/graph/model.h"
|
||||
#include "core/graph/graph_utils.h"
|
||||
#include "core/platform/env.h"
|
||||
#include "core/session/onnxruntime_cxx_api.h"
|
||||
#include "core/optimizer/reshape_fusion.h"
|
||||
#include "migraphx_inc.h"
|
||||
|
|
@ -108,6 +109,15 @@ MIGraphXExecutionProvider::MIGraphXExecutionProvider(const MIGraphXExecutionProv
|
|||
}
|
||||
|
||||
t_ = migraphx::target(info.target_device.c_str());
|
||||
|
||||
// Get environment variables
|
||||
const Env& env_instance = Env::Default();
|
||||
|
||||
// whether fp16 is enable
|
||||
const std::string fp16_enable_env = env_instance.GetEnvironmentVar(migraphx_env_vars::kFP16Enable);
|
||||
if (!fp16_enable_env.empty()) {
|
||||
fp16_enable_ = (std::stoi(fp16_enable_env) == 0 ? false : true);
|
||||
}
|
||||
}
|
||||
|
||||
AllocatorPtr MIGraphXExecutionProvider::GetAllocator(int id, OrtMemType mem_type) const {
|
||||
|
|
@ -197,33 +207,66 @@ static bool can_eval_concat(const Node* concat, const InitializedTensorSet& init
|
|||
{
|
||||
if (concat == nullptr) return true;
|
||||
const auto concat_args = concat->InputDefs();
|
||||
if (concat_args.size() != 3)
|
||||
{
|
||||
return false;
|
||||
}
|
||||
|
||||
auto arg_0 = concat_args[0];
|
||||
bool b_found = (initializers.find(arg_0->Name()) != initializers.end());
|
||||
auto arg_2 = concat_args[2];
|
||||
b_found &= (initializers.find(arg_2->Name()) != initializers.end());
|
||||
if (b_found)
|
||||
// scenario 1
|
||||
if (concat_args.size() == 1)
|
||||
{
|
||||
std::vector<graph_utils::EdgeEndToMatch> parent_path{
|
||||
{0, 1, "Unsqueeze", {1, 11}, kOnnxDomain},
|
||||
{0, 0, "Gather", {1, 11}, kOnnxDomain},
|
||||
{0, 0, "Shape", {1}, kOnnxDomain}};
|
||||
std::vector<graph_utils::EdgeEndToMatch> parent_path_1{
|
||||
{0, 0, "Unsqueeze", {1, 11}, kOnnxDomain},
|
||||
{0, 0, "Mul", {1, 6, 7}, kOnnxDomain},
|
||||
{0, 0, "Gather", {1, 11}, kOnnxDomain},
|
||||
{0, 0, "Shape", {1}, kOnnxDomain}
|
||||
};
|
||||
std::vector<const Node::EdgeEnd*> edges;
|
||||
b_found = graph_utils::FindPath(*concat, true, parent_path, edges, logger);
|
||||
bool b_found = graph_utils::FindPath(*concat, true, parent_path_1, edges, logger);
|
||||
if (b_found)
|
||||
{
|
||||
const Node& gather = edges[1]->GetNode();
|
||||
const auto* arg_index = gather.InputDefs()[1];
|
||||
if (initializers.find(arg_index->Name()) != initializers.end())
|
||||
const Node& mul = edges[1]->GetNode();
|
||||
const auto* arg_1 = mul.InputDefs()[1];
|
||||
bool const_flag = (initializers.find(arg_1->Name()) != initializers.end());
|
||||
if (const_flag)
|
||||
{
|
||||
return true;
|
||||
const Node& gather = edges[2]->GetNode();
|
||||
const auto* arg_index = gather.InputDefs()[1];
|
||||
if (initializers.find(arg_index->Name()) != initializers.end())
|
||||
{
|
||||
return true;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
else if (concat_args.size() >= 2)
|
||||
{
|
||||
int arg_size = static_cast<int>(concat_args.size());
|
||||
for (int i = 0; i < arg_size; ++i)
|
||||
{
|
||||
auto arg = concat_args[i];
|
||||
// is not an initializer
|
||||
if (initializers.find(arg->Name()) == initializers.end())
|
||||
{
|
||||
// then check whether can do constant folding for it
|
||||
std::vector<graph_utils::EdgeEndToMatch> parent_path{
|
||||
{0, i, "Unsqueeze", {1, 11}, kOnnxDomain},
|
||||
{0, 0, "Gather", {1, 11}, kOnnxDomain},
|
||||
{0, 0, "Shape", {1}, kOnnxDomain}};
|
||||
std::vector<const Node::EdgeEnd*> edges;
|
||||
bool b_found = graph_utils::FindPath(*concat, true, parent_path, edges, logger);
|
||||
if (!b_found)
|
||||
{
|
||||
return false;
|
||||
}
|
||||
|
||||
const Node& gather = edges[1]->GetNode();
|
||||
const auto* arg_index = gather.InputDefs()[1];
|
||||
if (initializers.find(arg_index->Name()) == initializers.end())
|
||||
{
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
return false;
|
||||
}
|
||||
|
|
@ -256,6 +299,7 @@ static bool can_eval_cast(const Node* cast, const InitializedTensorSet& initiali
|
|||
const_flag &= (initializers.find(slice_args[i]->Name()) != initializers.end());
|
||||
}
|
||||
}
|
||||
|
||||
if (const_flag)
|
||||
{
|
||||
return true;
|
||||
|
|
@ -345,7 +389,14 @@ static bool can_eval_input_shape(const Node* node, const InitializedTensorSet& i
|
|||
static bool IsUnsupportedOpMode(const Node* node, const onnxruntime::GraphViewer& graph_viewer, const logging::Logger& logger) {
|
||||
const auto& optype = node->OpType();
|
||||
const auto& initializers = graph_viewer.GetAllInitializedTensors();
|
||||
if (optype == "AveragePool") {
|
||||
if (optype == "ArgMax" or optype == "ArgMin") {
|
||||
const auto& attributes = node->GetAttributes();
|
||||
// we do not support select_last_index = 1 for now
|
||||
const auto sli_attr = attributes.find("select_last_index");
|
||||
if (sli_attr != attributes.end() && sli_attr->second.i() != 0) {
|
||||
return true;
|
||||
}
|
||||
} else if (optype == "AveragePool") {
|
||||
// ceil_mode attribute is not supported in MIGraphX
|
||||
const auto& attributes = node->GetAttributes();
|
||||
const auto ceil_attr = attributes.find("ceil_mode");
|
||||
|
|
@ -353,95 +404,6 @@ static bool IsUnsupportedOpMode(const Node* node, const onnxruntime::GraphViewer
|
|||
if (ceil_attr != attributes.end() && ceil_attr->second.i() != 0) {
|
||||
return true;
|
||||
}
|
||||
|
||||
// input can only have 4 dims
|
||||
const auto input_shape = node->InputDefs()[0]->Shape();
|
||||
if (input_shape != nullptr and input_shape->dim_size() != 4)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
|
||||
// migraphx does not support count_include_pad to be 1
|
||||
const auto cip_attr = attributes.find("count_include_pad");
|
||||
if (cip_attr != attributes.end() && cip_attr->second.i() != 0)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
|
||||
const auto ap_attr = attributes.find("auto_pad");
|
||||
if (ap_attr != attributes.end())
|
||||
{
|
||||
// explicit pad should be symmetric in migraphx
|
||||
auto s_pad = ap_attr->second.s();
|
||||
auto pads_attr = attributes.find("pads");
|
||||
if (s_pad == "NOTSET")
|
||||
{
|
||||
if (pads_attr != attributes.end())
|
||||
{
|
||||
auto pads = pads_attr->second.ints();
|
||||
if (pads.size() != 4)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
|
||||
if ((pads[0] != pads[2]) || (pads[1] != pads[3]))
|
||||
{
|
||||
return true;
|
||||
}
|
||||
}
|
||||
}
|
||||
// either SAME_UPPER or SAME_LOWER
|
||||
else if (s_pad.find("SAME") != std::string::npos)
|
||||
{
|
||||
// pads cannot exist when auto_pad is same_upper or same_lower
|
||||
if (pads_attr != attributes.end())
|
||||
{
|
||||
return true;
|
||||
}
|
||||
|
||||
// compute the padding size to see whether they are symmetric
|
||||
std::vector<int> strides = {1, 1};
|
||||
auto stride_attr = attributes.find("strides");
|
||||
if (stride_attr != attributes.end())
|
||||
{
|
||||
auto attr_strides = stride_attr->second.ints();
|
||||
strides.clear();
|
||||
std::copy(attr_strides.begin(), attr_strides.end(), std::back_inserter(strides));
|
||||
}
|
||||
|
||||
std::vector<int> kernel_lens = {1, 1};
|
||||
auto kernel_attr = attributes.find("kernel_shape");
|
||||
if (kernel_attr != attributes.end())
|
||||
{
|
||||
auto attr_k = kernel_attr->second.ints();
|
||||
std::copy(attr_k.begin(), attr_k.end(), kernel_lens.begin());
|
||||
}
|
||||
|
||||
auto tensor_dims = input_shape->dim();
|
||||
std::vector<int> in_lens;
|
||||
std::transform(tensor_dims.begin(),
|
||||
tensor_dims.end(),
|
||||
std::back_inserter(in_lens),
|
||||
[&](auto&& d) -> std::size_t {
|
||||
if(d.has_dim_value())
|
||||
{
|
||||
return d.dim_value();
|
||||
}
|
||||
return 1;
|
||||
});
|
||||
|
||||
std::vector<int> out_lens(2);
|
||||
out_lens[0] = (in_lens[2] + strides[0] - 1) / strides[0];
|
||||
out_lens[1] = (in_lens[3] + strides[1] - 1) / strides[1];
|
||||
std::vector<int> explicit_pads(2);
|
||||
explicit_pads[0] = (out_lens[0] - 1) * strides[0] + kernel_lens[0] - in_lens[2];
|
||||
explicit_pads[1] = (out_lens[1] - 1) * strides[1] + kernel_lens[1] - in_lens[3];
|
||||
if ((explicit_pads[0] & 1) != 0 or (explicit_pads[1] & 1) != 0)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
}
|
||||
}
|
||||
} else if (optype == "BatchNormalization") {
|
||||
// input can only have 4 dims
|
||||
const auto input_shape = node->InputDefs()[0]->Shape();
|
||||
|
|
@ -475,7 +437,11 @@ static bool IsUnsupportedOpMode(const Node* node, const onnxruntime::GraphViewer
|
|||
return false;
|
||||
}
|
||||
const Node* shape_node = graph_utils::GetInputNode(*node, 0);
|
||||
if (shape_node and shape_node->OpType() == "Concat")
|
||||
if (shape_node and shape_node->OpType() == "Shape")
|
||||
{
|
||||
return false;
|
||||
}
|
||||
else if (shape_node and shape_node->OpType() == "Concat")
|
||||
{
|
||||
if (can_eval_concat(shape_node, initializers, logger))
|
||||
{
|
||||
|
|
@ -517,6 +483,26 @@ static bool IsUnsupportedOpMode(const Node* node, const onnxruntime::GraphViewer
|
|||
// MIGraphX only supports constant shape input values
|
||||
const auto& shape_input = node->InputDefs()[1];
|
||||
return !graph_viewer.IsConstantInitializer(shape_input->Name(), true);
|
||||
} else if (optype == "Pow") {
|
||||
// we do not have a implementation to support different types of
|
||||
// the input data
|
||||
const auto args = node->InputDefs();
|
||||
const auto& input1_type = args[0]->TypeAsProto();
|
||||
if (input1_type == nullptr)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
auto data_type1 = input1_type->tensor_type().elem_type();
|
||||
const auto& input2_type = args[1]->TypeAsProto();
|
||||
if (input2_type == nullptr)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
auto data_type2 = input2_type->tensor_type().elem_type();
|
||||
if (data_type1 != data_type2)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
} else if (optype == "MaxPool") {
|
||||
//MaxPool "indices" output is not currently supported.
|
||||
if (node->OutputDefs().size() > 1) {
|
||||
|
|
@ -548,9 +534,15 @@ static bool IsUnsupportedOpMode(const Node* node, const onnxruntime::GraphViewer
|
|||
return true;
|
||||
}
|
||||
|
||||
// input can only have 4 dims
|
||||
const auto input_shape = node->InputDefs()[0]->Shape();
|
||||
if (input_shape != nullptr and input_shape->dim_size() != 4)
|
||||
// do not support int8 and uint8 type
|
||||
const auto& input_type = node->InputDefs()[0]->TypeAsProto();
|
||||
if (input_type == nullptr)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
auto data_type = input_type->tensor_type().elem_type();
|
||||
if (data_type == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_INT8 or
|
||||
data_type == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_UINT8)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
|
|
@ -631,6 +623,15 @@ static bool IsUnsupportedOpMode(const Node* node, const onnxruntime::GraphViewer
|
|||
{
|
||||
return false;
|
||||
}
|
||||
|
||||
const Node* shape_node = graph_utils::GetInputNode(*node, 0);
|
||||
if (shape_node and shape_node->OpType() == "Concat")
|
||||
{
|
||||
if (can_eval_concat(shape_node, initializers, logger))
|
||||
{
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
return true;
|
||||
} else if (optype == "Slice") {
|
||||
|
|
@ -668,6 +669,33 @@ static bool IsUnsupportedOpMode(const Node* node, const onnxruntime::GraphViewer
|
|||
}
|
||||
}
|
||||
}
|
||||
else if (optype == "Split")
|
||||
{
|
||||
// cannot process input dim of 0 size
|
||||
const auto arg_s = node->InputDefs()[0]->Shape();
|
||||
if (arg_s != nullptr)
|
||||
{
|
||||
auto tensor_dims = arg_s->dim();
|
||||
std::vector<std::size_t> dims;
|
||||
std::transform(tensor_dims.begin(),
|
||||
tensor_dims.end(),
|
||||
std::back_inserter(dims),
|
||||
[&](auto&& d) -> std::size_t {
|
||||
if(d.has_dim_value())
|
||||
{
|
||||
return d.dim_value();
|
||||
}
|
||||
else
|
||||
{
|
||||
return 0;
|
||||
}
|
||||
});
|
||||
if (dims == std::vector<std::size_t>{0})
|
||||
{
|
||||
return true;
|
||||
}
|
||||
}
|
||||
}
|
||||
else if (optype == "Tile")
|
||||
{
|
||||
const auto& args = node->InputDefs();
|
||||
|
|
@ -745,10 +773,11 @@ GetUnsupportedNodeIndices(const GraphViewer& graph_viewer,
|
|||
static std::set<std::string> mgx_supported_ops = {"Abs", "Acos", "Acosh", "Add", "ArgMax", "ArgMin",
|
||||
"Asin", "Asinh", "Atan", "Atanh", "AveragePool", "BatchNormalization", "Cast", "Ceil", "Clip",
|
||||
"Concat", "Constant", "ConstantFill", "ConstantOfShape", "Conv", "Cos", "Cosh", "Div", "Dropout",
|
||||
"Elu", "Erf", "Exp", "Expand", "Flatten", "Floor", "GRU", "Gather", "Gemm", "GlobalAveragePool",
|
||||
"GlobalMaxPool", "Identity", "ImageScaler", "InstanceNormalization", "LRN", "LSTM", "LeakyRelu",
|
||||
"Log", "LogSoftmax", "MatMul", "Max", "MaxPool", "Min", "Mul", "OneHot", "Pad", "Pow", "PRelu",
|
||||
"RNN","Range", "Reciprocal", "ReduceL1", "ReduceL2", "ReduceLogSum", "ReduceLogSumExp", "ReduceMax",
|
||||
"Elu", "Erf", "Exp", "Expand", "Flatten", "Floor", "GRU", "Gather", "GatherElements", "Gemm",
|
||||
"GlobalAveragePool", "GlobalMaxPool", "Identity", "ImageScaler", "InstanceNormalization", "LRN",
|
||||
"LSTM", "LeakyRelu", "Log", "LogSoftmax", "MatMul", "Max", "MaxPool", "Min", "Mul", "Neg",
|
||||
"OneHot", "Pad", "Pow", "PRelu",
|
||||
"RNN", "Range", "Reciprocal", "ReduceL1", "ReduceL2", "ReduceLogSum", "ReduceLogSumExp", "ReduceMax",
|
||||
"ReduceMean", "ReduceMin", "ReduceProd", "ReduceSum", "ReduceSumSquare", "Relu", "Reshape",
|
||||
"Round", "Shape", "Sigmoid", "Sign", "Sin", "Sinh", "Slice", "Softmax", "Split", "Sqrt", "Squeeze",
|
||||
"Sub", "Sum", "Tan", "Tanh", "Tile", "Transpose", "Unsqueeze"};
|
||||
|
|
@ -890,7 +919,6 @@ static void GetInputsOutputsOfSubgraph(const GraphViewer& graph_viewer,
|
|||
std::vector<std::unique_ptr<ComputeCapability>>
|
||||
MIGraphXExecutionProvider::GetCapability(const onnxruntime::GraphViewer& graph_viewer,
|
||||
const std::vector<const KernelRegistry*>& /*kernel_registries*/) const {
|
||||
|
||||
std::vector<std::unique_ptr<ComputeCapability>> result;
|
||||
if (graph_viewer.IsSubgraph()) {
|
||||
return result;
|
||||
|
|
@ -942,7 +970,6 @@ MIGraphXExecutionProvider::GetCapability(const onnxruntime::GraphViewer& graph_v
|
|||
// Example weights, reshape shape etc.
|
||||
std::unordered_set<std::string> mgx_required_initializers;
|
||||
const auto unsupported_nodes = GetUnsupportedNodeIndices(graph_viewer, mgx_required_initializers, *GetLogger());
|
||||
|
||||
// Too many unsupported operators, fallback to run on CPU
|
||||
if (unsupported_nodes.size() >= 6)
|
||||
{
|
||||
|
|
@ -1086,7 +1113,7 @@ Status MIGraphXExecutionProvider::Compile(const std::vector<onnxruntime::Node*>&
|
|||
std::string onnx_string_buffer;
|
||||
model_proto.SerializeToString(&onnx_string_buffer);
|
||||
std::vector<std::string> input_names, output_names;
|
||||
no_input_shape |= get_input_output_names(onnx_string_buffer, input_names, output_names);
|
||||
no_input_shape = no_input_shape or get_input_output_names(onnx_string_buffer, input_names, output_names);
|
||||
|
||||
// by parsing the model_proto, create a program corresponding to
|
||||
// the input fused_node
|
||||
|
|
@ -1095,6 +1122,10 @@ Status MIGraphXExecutionProvider::Compile(const std::vector<onnxruntime::Node*>&
|
|||
if (!no_input_shape)
|
||||
{
|
||||
prog = migraphx::parse_onnx_buffer(onnx_string_buffer, options);
|
||||
if (fp16_enable_)
|
||||
{
|
||||
migraphx::quantize_fp16(prog);
|
||||
}
|
||||
prog.compile(t_);
|
||||
|
||||
auto prog_output_shapes = prog.get_output_shapes();
|
||||
|
|
@ -1116,7 +1147,7 @@ Status MIGraphXExecutionProvider::Compile(const std::vector<onnxruntime::Node*>&
|
|||
std::unique_ptr<MIGraphXFuncState> p = onnxruntime::make_unique<MIGraphXFuncState>();
|
||||
*p = {context->allocate_func, context->release_func, context->allocator_handle, map_progs_[context->node_name],
|
||||
map_onnx_string_[context->node_name], options, t_, map_input_index_[context->node_name], &mgx_mu_,
|
||||
map_no_input_shape_[context->node_name]};
|
||||
map_no_input_shape_[context->node_name], fp16_enable_};
|
||||
*state = p.release();
|
||||
return 0;
|
||||
};
|
||||
|
|
@ -1135,6 +1166,7 @@ Status MIGraphXExecutionProvider::Compile(const std::vector<onnxruntime::Node*>&
|
|||
std::string& onnx_string = mgx_state->onnx_string;
|
||||
migraphx::onnx_options& cmp_options = mgx_state->options;
|
||||
bool &no_input_shape = mgx_state->no_input_shape;
|
||||
bool fp16_enable = mgx_state->fp16_enable;
|
||||
|
||||
// mean no program at all, so need to get the input shape info
|
||||
// from input data
|
||||
|
|
@ -1196,6 +1228,11 @@ Status MIGraphXExecutionProvider::Compile(const std::vector<onnxruntime::Node*>&
|
|||
if (!input_shape_match)
|
||||
{
|
||||
prog = migraphx::parse_onnx_buffer(onnx_string, cmp_options);
|
||||
if (fp16_enable)
|
||||
{
|
||||
migraphx::quantize_fp16(prog);
|
||||
}
|
||||
|
||||
prog.compile(t);
|
||||
mgx_state->prog = prog;
|
||||
param_shapes = prog.get_parameter_shapes();
|
||||
|
|
|
|||
|
|
@ -10,6 +10,10 @@
|
|||
|
||||
namespace onnxruntime {
|
||||
|
||||
namespace migraphx_env_vars {
|
||||
static const std::string kFP16Enable = "ORT_MIGRAPHX_FP16_ENABLE";
|
||||
};
|
||||
|
||||
// Information needed to construct amdmigraphx execution providers.
|
||||
struct MIGraphXExecutionProviderInfo {
|
||||
std::string target_device;
|
||||
|
|
@ -28,6 +32,7 @@ struct MIGraphXFuncState {
|
|||
std::unordered_map<std::string, std::size_t> input_name_indexes;
|
||||
OrtMutex* mgx_mu_ptr = nullptr;
|
||||
bool no_input_shape = false;
|
||||
bool fp16_enable = false;
|
||||
};
|
||||
|
||||
// Logical device representation.
|
||||
|
|
@ -48,6 +53,7 @@ class MIGraphXExecutionProvider : public IExecutionProvider {
|
|||
AllocatorPtr GetAllocator(int id, OrtMemType mem_type) const override;
|
||||
|
||||
private:
|
||||
bool fp16_enable_ = false;
|
||||
int device_id_;
|
||||
migraphx::target t_;
|
||||
OrtMutex mgx_mu_;
|
||||
|
|
|
|||
|
|
@ -80,7 +80,10 @@ def create_backend_test(testname=None):
|
|||
'^test_dynamicquantizelinear_max_adjusted_expanded_cpu', '^test_dynamicquantizelinear_min_adjusted_cpu',
|
||||
'^test_dynamicquantizelinear_min_adjusted_expanded_cpu',
|
||||
'^test_range_float_type_positive_delta_expanded_cpu',
|
||||
'^test_range_int32_type_negative_delta_expanded_cpu', '^test_operator_symbolic_override_nested_cpu'
|
||||
'^test_range_int32_type_negative_delta_expanded_cpu',
|
||||
'^test_operator_symbolic_override_nested_cpu',
|
||||
'^test_negative_log_likelihood_loss',
|
||||
'^test_softmax_cross_entropy'
|
||||
]
|
||||
|
||||
filters = current_failing_tests + \
|
||||
|
|
|
|||
Loading…
Reference in a new issue