From bfc888613f4e831d29c8b0bc17182ae061712553 Mon Sep 17 00:00:00 2001 From: Shucai Xiao Date: Thu, 25 Jun 2020 21:22:57 -0500 Subject: [PATCH] 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 --- BUILD.md | 24 ++ README.md | 2 +- dockerfiles/Dockerfile.migraphx | 8 +- dockerfiles/README.md | 18 +- docs/Roadmap.md | 1 - .../MIGraphX-ExecutionProvider.md | 43 +++ .../core/providers/migraphx/hip_allocator.cc | 8 - .../core/providers/migraphx/hip_allocator.h | 22 +- .../migraphx/migraphx_execution_provider.cc | 277 ++++++++++-------- .../migraphx/migraphx_execution_provider.h | 6 + .../test/python/onnx_backend_test_series.py | 5 +- 11 files changed, 270 insertions(+), 144 deletions(-) create mode 100644 docs/execution_providers/MIGraphX-ExecutionProvider.md diff --git a/BUILD.md b/BUILD.md index be15e76b13..d085169875 100644 --- a/BUILD.md +++ b/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 --use_migraphx --migraphx_home +``` + +Dockerfile instructions are available [here](./dockerfiles#migraphx) + *** # Training diff --git a/README.md b/README.md index d4a7386099..784382c1f1 100644 --- a/README.md +++ b/README.md @@ -137,7 +137,7 @@ For production scenarios, it's strongly recommended to build only from an [offic |CPU|GPU|IoT/Edge/Mobile|Other| |---|---|---|---| -|
  • Default CPU - *MLAS (Microsoft Linear Algebra Subprograms) + Eigen*
  • [Intel DNNL](./docs/execution_providers/DNNL-ExecutionProvider.md)
  • [Intel nGraph](./docs/execution_providers/nGraph-ExecutionProvider.md)
  • Intel MKL-ML *(build option)*
|
  • NVIDIA CUDA
  • [NVIDIA TensorRT](./docs/execution_providers/TensorRT-ExecutionProvider.md)
  • [DirectML](./docs/execution_providers/DirectML-ExecutionProvider.md)
|
  • [Intel OpenVINO](./docs/execution_providers/OpenVINO-ExecutionProvider.md)
  • [ARM Compute Library](./docs/execution_providers/ACL-ExecutionProvider.md) (*preview*)
  • [Android Neural Networks API](./docs/execution_providers/NNAPI-ExecutionProvider.md) (*preview*)
|
  • [Nuphar Model Compiler](./docs/execution_providers/Nuphar-ExecutionProvider.md) - (*preview*)
  • [Rockchip NPU](./docs/execution_providers/RKNPU-ExecutionProvider.md) (*preview*)
  • [Xilinx Vitis-AI](./docs/execution_providers/Vitis-AI-ExecutionProvider.md) (*preview*)
| +|
  • Default CPU - *MLAS (Microsoft Linear Algebra Subprograms) + Eigen*
  • [Intel DNNL](./docs/execution_providers/DNNL-ExecutionProvider.md)
  • [Intel nGraph](./docs/execution_providers/nGraph-ExecutionProvider.md)
  • Intel MKL-ML *(build option)*
|
  • NVIDIA CUDA
  • [NVIDIA TensorRT](./docs/execution_providers/TensorRT-ExecutionProvider.md)
  • [DirectML](./docs/execution_providers/DirectML-ExecutionProvider.md)
  • [AMD MIGraphX](./docs/execution_providers/MIGraphX-ExecutionProvider.md)
|
  • [Intel OpenVINO](./docs/execution_providers/OpenVINO-ExecutionProvider.md)
  • [ARM Compute Library](./docs/execution_providers/ACL-ExecutionProvider.md) (*preview*)
  • [Android Neural Networks API](./docs/execution_providers/NNAPI-ExecutionProvider.md) (*preview*)
|
  • [Nuphar Model Compiler](./docs/execution_providers/Nuphar-ExecutionProvider.md) - (*preview*)
  • [Rockchip NPU](./docs/execution_providers/RKNPU-ExecutionProvider.md) (*preview*)
  • [Xilinx Vitis-AI](./docs/execution_providers/Vitis-AI-ExecutionProvider.md) (*preview*)
| * [Roadmap: Upcoming accelerators](./docs/Roadmap.md#accelerators-and-execution-providers) * [Extensibility: Add an execution provider](docs/AddingExecutionProvider.md) diff --git a/dockerfiles/Dockerfile.migraphx b/dockerfiles/Dockerfile.migraphx index fcf33bee30..efcabf1698 100644 --- a/dockerfiles/Dockerfile.migraphx +++ b/dockerfiles/Dockerfile.migraphx @@ -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 diff --git a/dockerfiles/README.md b/dockerfiles/README.md index 715aaadb6c..f09fde2592 100644 --- a/dockerfiles/README.md +++ b/dockerfiles/README.md @@ -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* diff --git a/docs/Roadmap.md b/docs/Roadmap.md index 82cf54179b..305fe883a7 100644 --- a/docs/Roadmap.md +++ b/docs/Roadmap.md @@ -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 diff --git a/docs/execution_providers/MIGraphX-ExecutionProvider.md b/docs/execution_providers/MIGraphX-ExecutionProvider.md new file mode 100644 index 0000000000..96e74a48a3 --- /dev/null +++ b/docs/execution_providers/MIGraphX-ExecutionProvider.md @@ -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 +(std::unique_ptr{new CLogSink{}}, + static_cast(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. + diff --git a/onnxruntime/core/providers/migraphx/hip_allocator.cc b/onnxruntime/core/providers/migraphx/hip_allocator.cc index bd2b4c785d..7c8d6418f4 100644 --- a/onnxruntime/core/providers/migraphx/hip_allocator.cc +++ b/onnxruntime/core/providers/migraphx/hip_allocator.cc @@ -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(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(GetGPUDataTransfer(session_state)); } diff --git a/onnxruntime/core/providers/migraphx/hip_allocator.h b/onnxruntime/core/providers/migraphx/hip_allocator.h index bfd651ae08..8f6ea6dea9 100644 --- a/onnxruntime/core/providers/migraphx/hip_allocator.h +++ b/onnxruntime/core/providers/migraphx/hip_allocator.h @@ -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 diff --git a/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc b/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc index 4fbac1d6fa..5bca88e446 100644 --- a/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc +++ b/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc @@ -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 parent_path{ - {0, 1, "Unsqueeze", {1, 11}, kOnnxDomain}, - {0, 0, "Gather", {1, 11}, kOnnxDomain}, - {0, 0, "Shape", {1}, kOnnxDomain}}; + std::vector 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 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(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 parent_path{ + {0, i, "Unsqueeze", {1, 11}, kOnnxDomain}, + {0, 0, "Gather", {1, 11}, kOnnxDomain}, + {0, 0, "Shape", {1}, kOnnxDomain}}; + std::vector 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 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 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 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 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 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 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{0}) + { + return true; + } + } + } else if (optype == "Tile") { const auto& args = node->InputDefs(); @@ -745,10 +773,11 @@ GetUnsupportedNodeIndices(const GraphViewer& graph_viewer, static std::set 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> MIGraphXExecutionProvider::GetCapability(const onnxruntime::GraphViewer& graph_viewer, const std::vector& /*kernel_registries*/) const { - std::vector> 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 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& std::string onnx_string_buffer; model_proto.SerializeToString(&onnx_string_buffer); std::vector 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& 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& std::unique_ptr p = onnxruntime::make_unique(); *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& 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& 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(); diff --git a/onnxruntime/core/providers/migraphx/migraphx_execution_provider.h b/onnxruntime/core/providers/migraphx/migraphx_execution_provider.h index ee9dd3324c..36fc87d922 100644 --- a/onnxruntime/core/providers/migraphx/migraphx_execution_provider.h +++ b/onnxruntime/core/providers/migraphx/migraphx_execution_provider.h @@ -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 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_; diff --git a/onnxruntime/test/python/onnx_backend_test_series.py b/onnxruntime/test/python/onnx_backend_test_series.py index 3e0323e26f..47639da7d6 100644 --- a/onnxruntime/test/python/onnx_backend_test_series.py +++ b/onnxruntime/test/python/onnx_backend_test_series.py @@ -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 + \