Merge branch 'master' into ywang2/disable_CPU_model_cache_for_fallback
This commit is contained in:
Wang, Yang
GitHub
commit
8df4b685cf
|
|
@ -6,7 +6,7 @@ Additional Resources
|
|||
|
||||
|
||||
.. meta::
|
||||
:description: Learn more about OpenVINO from benchmark results, case studies
|
||||
:description: Learn more about OpenVINO from benchmark results, case studies
|
||||
and lists of supported models, operations and devices.
|
||||
|
||||
.. toctree::
|
||||
|
|
@ -14,7 +14,7 @@ Additional Resources
|
|||
:hidden:
|
||||
|
||||
additional-resources/glossary
|
||||
additional-resources/legal-information
|
||||
Legal and Responsible AI Information <./additional-resources/legal-information>
|
||||
additional-resources/telemetry
|
||||
Case Studies <https://www.intel.com/openvino-success-stories>
|
||||
|
||||
|
|
@ -23,9 +23,9 @@ Additional Resources
|
|||
|
||||
:doc:`Glossary <additional-resources/glossary>` contains terms used in OpenVINO.
|
||||
|
||||
:doc:`Legal Information <additional-resources/legal-information>` has trademark information and other legal statements.
|
||||
:doc:`Legal and Responsible AI Information <additional-resources/legal-information>` provides trademark information and other legal statements.
|
||||
|
||||
:doc:`OpenVINO™ Telemetry <additional-resources/telemetry>` has detailed information on the telemetry data collection.
|
||||
:doc:`OpenVINO™ Telemetry <additional-resources/telemetry>` has detailed information on the telemetry data collection.
|
||||
|
||||
`Case Studies <https://www.intel.com/openvino-success-stories>`__ are articles about real-world examples of OpenVINO™ usage.
|
||||
|
||||
|
|
|
|||
|
|
@ -1,7 +1,7 @@
|
|||
.. {#openvino_docs_Legal_Information}
|
||||
|
||||
Legal Information
|
||||
=================
|
||||
Legal and Responsible AI Information
|
||||
=====================================
|
||||
|
||||
|
||||
.. meta::
|
||||
|
|
@ -46,4 +46,12 @@ Intel Global Human Right Principles
|
|||
Intel is committed to respecting human rights and avoiding causing or contributing to adverse
|
||||
impacts on human rights. See `Intel's Global Human Rights Principles <https://www.intel.com/content/dam/www/central-libraries/us/en/documents/policy-human-rights.pdf>`__.
|
||||
Intel's products and software are intended only to be used in applications that do not cause or
|
||||
contribute to adverse impacts on human rights.
|
||||
contribute to adverse impacts on human rights.
|
||||
|
||||
Model Card Statement
|
||||
###########################################################
|
||||
|
||||
We recommend that users, wherever you are sourcing the model from, should check for a model card,
|
||||
consult the model card for each model you access and use, and create one if you are developing
|
||||
or updating a model. A model card is a short document that provides key information to assess
|
||||
performance and validation and ensure appropriate use.
|
||||
|
|
@ -1,5 +1,3 @@
|
|||
.. {#system_requirements}
|
||||
|
||||
System Requirements
|
||||
===================
|
||||
|
||||
|
|
@ -30,6 +28,7 @@ CPU
|
|||
|
||||
.. tab-item:: Supported Operating Systems
|
||||
|
||||
* Ubuntu 24.04 long-term support (LTS), 64-bit (Kernel 6.8+)
|
||||
* Ubuntu 22.04 long-term support (LTS), 64-bit (Kernel 5.15+)
|
||||
* Ubuntu 20.04 long-term support (LTS), 64-bit (Kernel 5.15+)
|
||||
* Ubuntu 18.04 long-term support (LTS) with limitations, 64-bit (Kernel 5.4+)
|
||||
|
|
@ -59,6 +58,7 @@ GPU
|
|||
|
||||
.. tab-item:: Supported Operating Systems
|
||||
|
||||
* Ubuntu 24.04 long-term support (LTS), 64-bit
|
||||
* Ubuntu 22.04 long-term support (LTS), 64-bit
|
||||
* Ubuntu 20.04 long-term support (LTS), 64-bit
|
||||
* Windows 10, 64-bit
|
||||
|
|
@ -75,7 +75,7 @@ GPU
|
|||
for information about your processor.
|
||||
* While this release of OpenVINO supports Ubuntu 20.04, the driver stack
|
||||
for Intel discrete graphic cards does not fully support Ubuntu 20.04.
|
||||
We recommend using Ubuntu 22.04 when executing on discrete graphics.
|
||||
We recommend using Ubuntu 22.04 and later when executing on discrete graphics.
|
||||
* The following minimum (i.e., used for old hardware) OpenCL™ driver's versions
|
||||
were used during OpenVINO internal validation: 22.43 for Ubuntu 22.04, 21.48
|
||||
for Ubuntu 20.04 and 21.49 for Red Hat Enterprise Linux 8 (some hardware may require
|
||||
|
|
@ -88,6 +88,7 @@ Intel® Neural Processing Unit
|
|||
|
||||
.. tab-item:: Operating Systems for NPU
|
||||
|
||||
* Ubuntu 24.04 long-term support (LTS), 64-bit
|
||||
* Ubuntu 22.04 long-term support (LTS), 64-bit
|
||||
* Windows 11, 64-bit (22H2, 23H2)
|
||||
|
||||
|
|
@ -106,6 +107,7 @@ Operating systems and developer environment
|
|||
|
||||
.. tab-item:: Linux OS
|
||||
|
||||
* Ubuntu 24.04 with Linux kernel 6.8+
|
||||
* Ubuntu 22.04 with Linux kernel 5.15+
|
||||
* Ubuntu 20.04 with Linux kernel 5.15+
|
||||
* Red Hat Enterprise Linux 8 with Linux kernel 5.4
|
||||
|
|
|
|||
|
|
@ -16,7 +16,7 @@ GET STARTED
|
|||
Install OpenVINO <get-started/install-openvino>
|
||||
Additional Hardware Setup <get-started/configurations>
|
||||
Troubleshooting <get-started/troubleshooting-install-config>
|
||||
System Requirements <about-openvino/system-requirements>
|
||||
System Requirements <./about-openvino/release-notes-openvino/system-requirements>
|
||||
|
||||
|
||||
.. raw:: html
|
||||
|
|
|
|||
|
|
@ -32,12 +32,16 @@ Below are the instructions on how to install the OpenCL packages on supported Li
|
|||
|
||||
.. tab-set::
|
||||
|
||||
.. tab-item:: Ubuntu 22.04 LTS
|
||||
.. tab-item:: Ubuntu 22.04 LTS / Ubuntu 24.04 LTS
|
||||
:sync: ubuntu-22
|
||||
|
||||
Download and install the `deb` packages published `here <https://github.com/intel/compute-runtime/releases/latest>`__ and install the apt package `ocl-icd-libopencl1` with the OpenCl ICD loader.
|
||||
Download and install the `deb` packages published `here <https://github.com/intel/compute-runtime/releases/latest>`__
|
||||
and install the apt package `ocl-icd-libopencl1` with the OpenCl ICD loader.
|
||||
|
||||
Alternatively, you can add the apt repository by following the `installation guide <https://dgpu-docs.intel.com/driver/installation.html#ubuntu-install-steps>`__. Then install the `ocl-icd-libopencl1`, `intel-opencl-icd`, `intel-level-zero-gpu` and `level-zero` apt packages:
|
||||
Alternatively, you can add the apt repository by following the
|
||||
`installation guide <https://dgpu-docs.intel.com/driver/installation.html#ubuntu-install-steps>`__.
|
||||
Then install the `ocl-icd-libopencl1`, `intel-opencl-icd`, `intel-level-zero-gpu` and `level-zero`
|
||||
apt packages:
|
||||
|
||||
.. code-block:: sh
|
||||
|
||||
|
|
@ -119,7 +123,7 @@ Below are the required steps to make it work with OpenVINO:
|
|||
wsl --update
|
||||
wsl --shutdown
|
||||
|
||||
- When booting Ubuntu 20.04 or Ubuntu 22.04, install the same drivers as described above in the Linux section
|
||||
- When booting Ubuntu 20.04, 22.04, or 24.04 install the same drivers as described above in the Linux section
|
||||
|
||||
.. note::
|
||||
|
||||
|
|
@ -128,29 +132,6 @@ Below are the required steps to make it work with OpenVINO:
|
|||
Additional Resources
|
||||
####################
|
||||
|
||||
.. The following Intel® Graphics Driver versions were used during OpenVINO's internal validation:
|
||||
|
||||
.. <The table below is out of date and we do not have an updated list of drivers used for validation as of 2024.0 release date.>
|
||||
.. <The table will be updated when an updated list of drivers is available>
|
||||
|
||||
.. +------------------+-------------------------------------------------------------------------------------------+
|
||||
.. | Operation System | Driver version |
|
||||
.. +==================+===========================================================================================+
|
||||
.. | Ubuntu 22.04 | `22.43.24595.30 <https://github.com/intel/compute-runtime/releases/tag/22.43.24595.30>`__ |
|
||||
.. +------------------+-------------------------------------------------------------------------------------------+
|
||||
.. | Ubuntu 20.04 | `22.35.24055 <https://github.com/intel/compute-runtime/releases/tag/22.35.24055>`__ |
|
||||
.. +------------------+-------------------------------------------------------------------------------------------+
|
||||
.. | Ubuntu 18.04 | `21.38.21026 <https://github.com/intel/compute-runtime/releases/tag/21.38.21026>`__ |
|
||||
.. +------------------+-------------------------------------------------------------------------------------------+
|
||||
.. | CentOS 7 | `19.41.14441 <https://github.com/intel/compute-runtime/releases/tag/19.41.14441>`__ |
|
||||
.. +------------------+-------------------------------------------------------------------------------------------+
|
||||
.. | RHEL 8 | `22.28.23726 <https://github.com/intel/compute-runtime/releases/tag/22.28.23726>`__ |
|
||||
.. +------------------+-------------------------------------------------------------------------------------------+
|
||||
|
||||
|
||||
.. What’s Next?
|
||||
.. ############
|
||||
|
||||
* :doc:`GPU Device <../../openvino-workflow/running-inference/inference-devices-and-modes/gpu-device>`
|
||||
* :doc:`Install Intel® Distribution of OpenVINO™ toolkit from a Docker Image <../install-openvino/install-openvino-archive-linux>`
|
||||
* `Docker CI framework for Intel® Distribution of OpenVINO™ toolkit <https://github.com/openvinotoolkit/docker_ci/blob/master/README.md>`__
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ Configurations for Intel® NPU with OpenVINO™
|
|||
===============================================
|
||||
|
||||
.. meta::
|
||||
:description: Learn how to provide additional configuration for Intel®
|
||||
:description: Learn how to provide additional configuration for Intel®
|
||||
NPU to work with the OpenVINO™ toolkit on your system.
|
||||
|
||||
|
||||
|
|
@ -19,27 +19,27 @@ Make sure you use the most recent supported driver for your hardware setup.
|
|||
The driver is maintained as open source and may be found in the following repository,
|
||||
together with comprehensive information on installation and system requirements:
|
||||
`github.com/intel/linux-npu-driver <https://github.com/intel/linux-npu-driver>`__
|
||||
|
||||
|
||||
It is recommended to check for the latest version of the driver.
|
||||
|
||||
Make sure you use a supported OS version, as well as install make, gcc,
|
||||
and Linux kernel headers. To check the NPU state, use the ``dmesg``
|
||||
command in the console. A successful boot-up of the NPU should give you
|
||||
a message like this one:
|
||||
|
||||
|
||||
``[ 797.193201] [drm] Initialized intel_vpu 0.<version number> for 0000:00:0b.0 on minor 0``
|
||||
|
||||
The current requirement for inference on NPU is Ubuntu 22.04 with the kernel
|
||||
version of 6.6 or higher.
|
||||
The current requirement for inference on NPU is the minimum of Ubuntu 22.04, kernel
|
||||
version of 6.6.
|
||||
|
||||
.. tab-item:: Windows
|
||||
|
||||
The Intel® NPU driver for Windows is available through Windows Update but
|
||||
it may also be installed manually by downloading the
|
||||
`NPU driver package <https://www.intel.com/content/www/us/en/download/794734/intel-npu-driver-windows.html>`__ and following the
|
||||
it may also be installed manually by downloading the
|
||||
`NPU driver package <https://www.intel.com/content/www/us/en/download/794734/intel-npu-driver-windows.html>`__ and following the
|
||||
`Windows driver installation guide <https://support.microsoft.com/en-us/windows/update-drivers-manually-in-windows-ec62f46c-ff14-c91d-eead-d7126dc1f7b6>`__.
|
||||
|
||||
If a driver has already been installed you should be able to find
|
||||
'Intel(R) NPU Accelerator' in Windows Device Manager. If you
|
||||
If a driver has already been installed you should be able to find
|
||||
'Intel(R) NPU Accelerator' in Windows Device Manager. If you
|
||||
cannot find such a device, the NPU is most likely listed in "Other devices"
|
||||
as "Multimedia Video Controller."
|
||||
|
|
|
|||
|
|
@ -1,6 +1,4 @@
|
|||
.. {#openvino_docs_install_guides_overview}
|
||||
|
||||
Install OpenVINO™ 2024.0
|
||||
Install OpenVINO™ 2024.2
|
||||
==========================
|
||||
|
||||
|
||||
|
|
@ -36,17 +34,17 @@ Install OpenVINO™ 2024.0
|
|||
|
||||
.. tip::
|
||||
|
||||
OpenVINO 2024.0, described here, is not a Long-Term-Support version!
|
||||
OpenVINO 2024.2, described here, is not a Long-Term-Support version!
|
||||
All currently supported versions are:
|
||||
|
||||
* 2024.0 (development)
|
||||
* 2024.2 (development)
|
||||
* 2023.3 (LTS)
|
||||
* 2022.3 (LTS)
|
||||
|
||||
Moreover, different OpenVINO distributions may support slightly different sets of features.
|
||||
Read installation guides for particular distributions for more details.
|
||||
|
||||
.. dropdown:: Distribution Comparison for OpenVINO 2024.0
|
||||
.. dropdown:: Distribution Comparison for OpenVINO 2024.2
|
||||
|
||||
=============== ========== ====== =============== ======== ============ ========== ========== ==========
|
||||
Device Archives PyPI APT/YUM/ZYPPER Conda Homebrew vcpkg Conan npm
|
||||
|
|
@ -56,7 +54,7 @@ Install OpenVINO™ 2024.0
|
|||
NPU V\* V\* V\ * n/a n/a n/a n/a V\*
|
||||
=============== ========== ====== =============== ======== ============ ========== ========== ==========
|
||||
|
||||
| \* **Of the Linux systems, only Ubuntu 22.04 includes drivers for NPU device.**
|
||||
| \* **Of the Linux systems, versions 22.04 and 24.04 include drivers for NPU.**
|
||||
| **For Windows, CPU inference on ARM64 is not supported.**
|
||||
|
||||
| **Build OpenVINO from source**
|
||||
|
|
|
|||
|
|
@ -73,6 +73,13 @@ Step 1: Set Up the OpenVINO Toolkit APT Repository
|
|||
|
||||
.. tab-set::
|
||||
|
||||
.. tab-item:: Ubuntu 24
|
||||
:sync: ubuntu-24
|
||||
|
||||
.. code-block:: sh
|
||||
|
||||
echo "deb https://apt.repos.intel.com/openvino/2024 ubuntu24 main" | sudo tee /etc/apt/sources.list.d/intel-openvino-2024.list
|
||||
|
||||
.. tab-item:: Ubuntu 22
|
||||
:sync: ubuntu-22
|
||||
|
||||
|
|
|
|||
|
|
@ -1,5 +1,3 @@
|
|||
.. {#openvino_docs_install_guides_installing_openvino_from_archive_linux}
|
||||
|
||||
Install OpenVINO™ Runtime on Linux from an Archive File
|
||||
=========================================================
|
||||
|
||||
|
|
@ -30,6 +28,7 @@ Install OpenVINO™ Runtime on Linux from an Archive File
|
|||
Ubuntu18 x86_64 V V n/a
|
||||
Ubuntu20 x86_64 V V V
|
||||
Ubuntu22 x86_64 V V V
|
||||
Ubuntu24 x86_64 V V V
|
||||
RHEL8 x86_64 V V n/a
|
||||
=================== ===== ===== =====
|
||||
|
||||
|
|
@ -130,6 +129,16 @@ Step 1: Download and Install the OpenVINO Core Components
|
|||
|
||||
.. tab-set::
|
||||
|
||||
.. tab-item:: Ubuntu 24.04
|
||||
:sync: ubuntu-24
|
||||
|
||||
.. code-block:: sh
|
||||
|
||||
|
||||
curl -L https://storage.openvinotoolkit.org/repositories/openvino/packages/2024.1/linux/l_openvino_toolkit_ubuntu22_2024.1.0.15008.f4afc983258_x86_64.tgz --output openvino_2024.1.0.tgz
|
||||
tar -xf openvino_2024.1.0.tgz
|
||||
sudo mv l_openvino_toolkit_ubuntu24_2024.1.0.15008.f4afc983258_x86_64 /opt/intel/openvino_2024.1.0
|
||||
|
||||
.. tab-item:: Ubuntu 22.04
|
||||
:sync: ubuntu-22
|
||||
|
||||
|
|
|
|||
|
|
@ -10,14 +10,14 @@ Troubleshooting Guide for OpenVINO™ Installation & Configuration
|
|||
of OpenVINO™ on your system.
|
||||
|
||||
|
||||
| This guide provides general troubleshooting steps and solutions to possible issues that
|
||||
may be encountered while installing and configuring OpenVINO™. For a comprehensive
|
||||
database of support topics on OpenVINO, go to:
|
||||
| This article provides general troubleshooting steps and solutions to possible issues that you
|
||||
may face while installing and configuring OpenVINO™. For a comprehensive database of support
|
||||
topics on OpenVINO, go to:
|
||||
| `Support for OpenVINO™ toolkit <https://www.intel.com/content/www/us/en/support/products/96066/software/development-software/openvino-toolkit.html>`__
|
||||
|
||||
|
||||
|
||||
.. dropdown:: Errors with Installing via PIP for Users in China
|
||||
.. dropdown:: PIP for Users in China gives errors
|
||||
|
||||
Users in China might encounter errors while downloading sources via PIP during OpenVINO™
|
||||
installation. To resolve the issues, try adding the download source using the ``-i``
|
||||
|
|
@ -34,16 +34,18 @@ Troubleshooting Guide for OpenVINO™ Installation & Configuration
|
|||
|
||||
pip install openvino-dev[tensorflow2] -i https://mirrors.aliyun.com/pypi/simple/
|
||||
|
||||
.. dropdown:: ImportError: cannot import name 'Core' from 'openvino'
|
||||
|
||||
.. dropdown:: Issues with Installing OpenVINO on Linux from Docker
|
||||
This error may appear on systems lacking C++ components. Since it is almost exclusively a
|
||||
Windows case, installing `Microsoft Visual C++ Redistributable [vc_redist.x64] <https://aka.ms/vs/17/release/vc_redist.x64.exe>`__
|
||||
package may fix it. For more information on dependencies, check
|
||||
:doc:`System Requirements <../about-openvino/release-notes-openvino/system-requirements>` and
|
||||
:doc:`Additional Hardware Configurations <./configurations>`
|
||||
|
||||
.. _proxy-issues:
|
||||
.. dropdown:: Proxy issues installing OpenVINO on Linux from Docker
|
||||
|
||||
Proxy Issues
|
||||
++++++++++++
|
||||
|
||||
If you meet proxy issues during the installation with Docker, you need set up proxy settings
|
||||
for Docker. See the `Docker guide <https://docs.docker.com/network/proxy/#set-proxy-using-the-cli>`__
|
||||
If you face proxy issues during installation with Docker, you may need to set up proxy
|
||||
settings for it. See the `Docker guide <https://docs.docker.com/network/proxy/#set-proxy-using-the-cli>`__
|
||||
for more details.
|
||||
|
||||
.. dropdown:: Check the version of OpenVINO Runtime
|
||||
|
|
@ -81,11 +83,11 @@ Troubleshooting Guide for OpenVINO™ Installation & Configuration
|
|||
|
||||
.. dropdown:: Check if environment variables are set correctly
|
||||
|
||||
- For Python developers, if you previously installed OpenVINO using the archive file,
|
||||
* For Python developers, if you previously installed OpenVINO using the archive file,
|
||||
and are now installing OpenVINO using PIP, remove all the PATH settings and the lines with
|
||||
``setupvars`` from ``.bashrc``. Note that if you installed OpenVINO with PIP in a virtual
|
||||
environment, you don't need to set any environment variables.
|
||||
- If you have installed OpenVINO before, you probably have added ``setupvars`` to your
|
||||
* If you have installed OpenVINO before, you probably have added ``setupvars`` to your
|
||||
``PATH /.bashrc`` or Windows environment variables. After restarting your environment,
|
||||
you should see an information similar to the following:
|
||||
|
||||
|
|
@ -93,10 +95,10 @@ Troubleshooting Guide for OpenVINO™ Installation & Configuration
|
|||
|
||||
[setupvars.sh] OpenVINO™ environment initialized
|
||||
|
||||
- If you don't see the information above, your PATH variables may be configured incorrectly.
|
||||
* If you don't see the information above, your PATH variables may be configured incorrectly.
|
||||
Check if you have typed the correct <INSTALL_DIR> or you are trying to activate in the
|
||||
correct directory.
|
||||
- If you added it to a ``.bashrc`` file, make sure that the command is correctly written and
|
||||
* If you added it to a ``.bashrc`` file, make sure that the command is correctly written and
|
||||
the file is found in the ``~/.bashrc`` folder.
|
||||
|
||||
.. dropdown:: Verify that OpenVINO is correctly installed
|
||||
|
|
|
|||
|
|
@ -1,5 +1,3 @@
|
|||
.. {#openvino_deployment_guide}
|
||||
|
||||
Deploy Locally
|
||||
==============
|
||||
|
||||
|
|
@ -43,11 +41,11 @@ The table below shows which distribution type can be used for what target operat
|
|||
* - Distribution type
|
||||
- Operating systems
|
||||
* - Debian packages
|
||||
- Ubuntu 18.04 long-term support (LTS), 64-bit; Ubuntu 20.04 long-term support (LTS), 64-bit
|
||||
- Ubuntu 18.04, 20.04, 22.04, 24.04 (64-bit)
|
||||
* - RPM packages
|
||||
- Red Hat Enterprise Linux 8, 64-bit
|
||||
* - Docker images
|
||||
- Ubuntu 22.04 long-term support (LTS), 64-bit; Ubuntu 20.04 long-term support (LTS), 64-bit; Red Hat Enterprise Linux 8, 64-bit
|
||||
- Ubuntu 20.04, 22.04, 24.04 (64-bit); Red Hat Enterprise Linux 8, 64-bit
|
||||
* - PyPI (PIP package manager)
|
||||
- See https://pypi.org/project/openvino
|
||||
* - :doc:`Libraries for Local Distribution <deployment-locally/local-distribution-libraries>`
|
||||
|
|
|
|||
|
|
@ -1,5 +1,3 @@
|
|||
.. {#openvino_docs_OV_UG_supported_plugins_NPU}
|
||||
|
||||
NPU Device
|
||||
==========
|
||||
|
||||
|
|
|
|||
|
|
@ -7,7 +7,7 @@ The software was validated on:
|
|||
|
||||
> **NOTE**: To build on CentOS 7 (64-bit), refer to [Building OpenVINO on CentOS 7 Guide](https://github.com/openvinotoolkit/openvino/wiki/Building-OpenVINO-on-CentOS-7-Guide)
|
||||
|
||||
## Software requirements
|
||||
## Software requirements
|
||||
|
||||
- [CMake](https://cmake.org/download/) 3.13 or higher
|
||||
- GCC 7.5 or higher to build OpenVINO Runtime
|
||||
|
|
@ -37,9 +37,6 @@ The software was validated on:
|
|||
2. Install build dependencies using the `install_build_dependencies.sh` script in the
|
||||
project root folder.
|
||||
```sh
|
||||
chmod +x install_build_dependencies.sh
|
||||
```
|
||||
```sh
|
||||
sudo ./install_build_dependencies.sh
|
||||
```
|
||||
|
||||
|
|
@ -79,7 +76,7 @@ You can use the following additional build options:
|
|||
```sh
|
||||
pip install -r <openvino source tree>/src/bindings/python/wheel/requirements-dev.txt
|
||||
```
|
||||
3. After the build process finishes, export the newly built Python libraries to the user environment variables:
|
||||
3. After the build process finishes, export the newly built Python libraries to the user environment variables:
|
||||
```
|
||||
export PYTHONPATH=<openvino_repo>/bin/intel64/Release/python:$PYTHONPATH
|
||||
export LD_LIBRARY_PATH=<openvino_repo>/bin/intel64/Release:$LD_LIBRARY_PATH
|
||||
|
|
|
|||
|
|
@ -36,6 +36,21 @@ git clone --recurse-submodules --single-branch --branch=master https://github.co
|
|||
.. && cmake --build . --parallel
|
||||
```
|
||||
|
||||
> **NOTE**: The build command may fail due to insufficient RAM. To fix this issue, you can increase the swap size:
|
||||
1. Deactivate the current swap:
|
||||
```bash
|
||||
sudo dphys-swapfile swapoff
|
||||
```
|
||||
2. Modify the swap size by setting `CONF_SWAPSIZE=8192` in `/etc/dphys-swapfile`.
|
||||
3. Recreate the swap file:
|
||||
```bash
|
||||
sudo dphys-swapfile setup
|
||||
```
|
||||
3. Start swap:
|
||||
```bash
|
||||
sudo dphys-swapfile swapon
|
||||
```
|
||||
|
||||
## Additional Build Options
|
||||
|
||||
- To build Python API, install `libpython3-dev:armhf` and `python3-pip`
|
||||
|
|
|
|||
|
|
@ -2,7 +2,7 @@
|
|||
|
||||
OpenVINO components provides different debug capabilities, to get more information please read:
|
||||
|
||||
* [OpenVINO Model Debug Capabilities](https://docs.openvino.ai/2024/openvino_docs_OV_UG_Model_Representation.html#model-debugging-capabilities)
|
||||
* [OpenVINO Model Debug Capabilities](https://docs.openvino.ai/2024/openvino-workflow/running-inference/integrate-openvino-with-your-application/model-representation.html#model-debugging-capabilities)
|
||||
* [OpenVINO Pass Manager Debug Capabilities](#todo)
|
||||
|
||||
## See also
|
||||
|
|
|
|||
|
|
@ -107,7 +107,9 @@ def test_custom_add_model():
|
|||
|
||||
def test_custom_op():
|
||||
model = create_snake_model()
|
||||
compiled_model = compile_model(model)
|
||||
# todo: CVS-141744
|
||||
# it hangs with AUTO plugin, but works well with CPU
|
||||
compiled_model = compile_model(model, "CPU")
|
||||
|
||||
assert isinstance(compiled_model, CompiledModel)
|
||||
request = compiled_model.create_infer_request()
|
||||
|
|
|
|||
|
|
@ -23,6 +23,5 @@ class PrevSequenceLengthPattern;
|
|||
class ov::pass::PrevSequenceLengthPattern : public ov::pass::MatcherPass {
|
||||
public:
|
||||
OPENVINO_RTTI("PrevSequenceLengthPattern", "0");
|
||||
explicit PrevSequenceLengthPattern(const std::shared_ptr<ov::op::v1::Subtract>& prev_max_seq_len,
|
||||
std::shared_ptr<ov::Node>);
|
||||
explicit PrevSequenceLengthPattern(std::shared_ptr<ov::Node> prev_max_seq_len, std::shared_ptr<ov::Node> batch_dim);
|
||||
};
|
||||
|
|
@ -14,9 +14,8 @@
|
|||
|
||||
using namespace ov::op;
|
||||
|
||||
ov::pass::PrevSequenceLengthPattern::PrevSequenceLengthPattern(
|
||||
const std::shared_ptr<ov::op::v1::Subtract>& prev_max_seq_len,
|
||||
std::shared_ptr<ov::Node> batch_dim) {
|
||||
ov::pass::PrevSequenceLengthPattern::PrevSequenceLengthPattern(std::shared_ptr<ov::Node> prev_max_seq_len,
|
||||
std::shared_ptr<ov::Node> batch_dim) {
|
||||
MATCHER_SCOPE(PrevSequenceLengthPattern);
|
||||
// The transformation addresses two cases that look similar: (1) previous sequence length, (2) batch size in
|
||||
// kv-cache state In first case it should replace it by prev_max_seq_len. For the second case, connect to batch_dim.
|
||||
|
|
@ -39,30 +38,23 @@ ov::pass::PrevSequenceLengthPattern::PrevSequenceLengthPattern(
|
|||
auto axis = gather_index->cast_vector<int64_t>().at(0);
|
||||
auto kv_init_shape = pattern_map.at(kv_past).get_node()->get_input_partial_shape(0);
|
||||
auto target_type = gather->get_output_element_type(0);
|
||||
std::shared_ptr<ov::Node> replacement;
|
||||
if (kv_init_shape[axis].is_static() && kv_init_shape[axis].get_length() == 0) {
|
||||
// this is a sequence dimension based on how the initialization expression is build for stateful models
|
||||
std::shared_ptr<ov::Node> replacement;
|
||||
if (prev_max_seq_len->get_output_element_type(0) != target_type) {
|
||||
replacement = std::make_shared<v0::Convert>(prev_max_seq_len, target_type);
|
||||
} else {
|
||||
replacement = prev_max_seq_len;
|
||||
}
|
||||
replace_node(
|
||||
gather,
|
||||
std::make_shared<v1::Reshape>(replacement, v0::Constant::create(element::i64, Shape{1}, {1}), false));
|
||||
return true;
|
||||
} else { // assumption that any other axis should point to batch dimension, precise reasoning is too complex
|
||||
// (TODO)
|
||||
// this is a batch dimension
|
||||
std::shared_ptr<ov::Node> replacement;
|
||||
if (batch_dim->get_output_element_type(0) != target_type) {
|
||||
replacement = std::make_shared<v0::Convert>(batch_dim, target_type);
|
||||
} else {
|
||||
replacement = batch_dim;
|
||||
}
|
||||
replace_node(gather, replacement);
|
||||
return true;
|
||||
replacement = prev_max_seq_len;
|
||||
} else {
|
||||
// assumption that any other axis should point to batch dimension, precise reasoning is too complex
|
||||
// TODO: provide more reliable check
|
||||
replacement = batch_dim;
|
||||
}
|
||||
if (replacement->get_output_element_type(0) != target_type) {
|
||||
replacement = std::make_shared<v0::Convert>(replacement, target_type);
|
||||
}
|
||||
auto required_shape = gather->get_output_partial_shape(0);
|
||||
if (replacement->get_output_partial_shape(0) != required_shape && required_shape.rank().is_static()) {
|
||||
replacement = op::util::reshapeTo(replacement, Shape(required_shape.rank().get_length(), 1));
|
||||
}
|
||||
replace_node(gather, replacement);
|
||||
return true;
|
||||
};
|
||||
|
||||
auto m = std::make_shared<ov::pass::pattern::Matcher>(seq, matcher_name);
|
||||
|
|
|
|||
|
|
@ -7,6 +7,7 @@
|
|||
#include "openvino/cc/pass/itt.hpp"
|
||||
#include "openvino/op/concat.hpp"
|
||||
#include "openvino/op/gather.hpp"
|
||||
#include "openvino/op/reshape.hpp"
|
||||
#include "openvino/op/shape_of.hpp"
|
||||
#include "openvino/pass/pattern/op/wrap_type.hpp"
|
||||
#include "transformations/utils/utils.hpp"
|
||||
|
|
@ -29,11 +30,13 @@ ov::pass::TotalSequenceLengthPattern::TotalSequenceLengthPattern(
|
|||
// use symbolic infra or look at the constant input
|
||||
auto gather = m.get_match_root();
|
||||
auto target_type = gather->get_output_element_type(0);
|
||||
std::shared_ptr<Node> replacement;
|
||||
if (max_context_len->get_output_element_type(0) != target_type) {
|
||||
replacement = std::make_shared<v0::Convert>(max_context_len, target_type);
|
||||
} else {
|
||||
replacement = max_context_len;
|
||||
std::shared_ptr<Node> replacement = max_context_len;
|
||||
if (replacement->get_output_element_type(0) != target_type) {
|
||||
replacement = std::make_shared<v0::Convert>(replacement, target_type);
|
||||
}
|
||||
auto required_shape = gather->get_output_partial_shape(0);
|
||||
if (replacement->get_output_partial_shape(0) != required_shape && required_shape.rank().is_static()) {
|
||||
replacement = op::util::reshapeTo(replacement, Shape(required_shape.rank().get_length(), 1));
|
||||
}
|
||||
replace_node(gather, replacement);
|
||||
return true;
|
||||
|
|
|
|||
|
|
@ -36,7 +36,12 @@ bool relax_hc_reshape_followed_by_matmul(const ov::pass::pattern::PatternValueMa
|
|||
// avoiding loop creation
|
||||
return false;
|
||||
|
||||
const auto idx = reshape_is_A_input ? (matmul->get_transpose_b() ? -1 : -2) : (matmul->get_transpose_a() ? -2 : -1);
|
||||
bool is_1d = ov::pass::pattern::rank_equals(1)(shape_source);
|
||||
int64_t idx = -1;
|
||||
if (!is_1d) {
|
||||
idx = reshape_is_A_input ? (matmul->get_transpose_b() ? -1 : -2) : (matmul->get_transpose_a() ? -2 : -1);
|
||||
}
|
||||
|
||||
const auto in_C_0 = std::make_shared<ov::op::v3::ShapeOf>(shape_source);
|
||||
const auto in_C_1 = ov::op::v0::Constant::create(ov::element::i64, {1}, {idx});
|
||||
const auto in_C_2 = ov::op::v0::Constant::create(ov::element::i64, {}, {0});
|
||||
|
|
@ -53,7 +58,6 @@ bool relax_hc_reshape_followed_by_matmul(const ov::pass::pattern::PatternValueMa
|
|||
|
||||
auto reshape_input = pattern_to_output.at(reshape_label).get_node_shared_ptr()->input(1);
|
||||
reshape_input.replace_source_output(new_reshape_pattern);
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
|
|
@ -61,7 +65,7 @@ bool relax_hc_reshape_followed_by_matmul(const ov::pass::pattern::PatternValueMa
|
|||
|
||||
ov::pass::ReshapeAMatMul::ReshapeAMatMul() {
|
||||
MATCHER_SCOPE(ReshapeAMatMul);
|
||||
auto other_input_label = pattern::any_input();
|
||||
auto other_input_label = pattern::any_input(ov::pass::pattern::has_static_rank());
|
||||
auto reshape_input_label = pattern::any_input();
|
||||
auto reshape_pattern_label = pattern::any_input();
|
||||
auto reshape_predicate = [](ov::Output<ov::Node> output) -> bool {
|
||||
|
|
@ -86,7 +90,7 @@ ov::pass::ReshapeAMatMul::ReshapeAMatMul() {
|
|||
|
||||
ov::pass::ReshapeBMatMul::ReshapeBMatMul() {
|
||||
MATCHER_SCOPE(ReshapeBMatMul);
|
||||
auto other_input_label = pattern::any_input();
|
||||
auto other_input_label = pattern::any_input(ov::pass::pattern::has_static_rank());
|
||||
auto reshape_input_label = pattern::any_input();
|
||||
auto reshape_pattern_label = pattern::any_input();
|
||||
auto reshape_predicate = [](ov::Output<ov::Node> output) -> bool {
|
||||
|
|
|
|||
|
|
@ -50,6 +50,10 @@ public:
|
|||
return m_causal;
|
||||
}
|
||||
|
||||
void set_causal(bool causal) {
|
||||
m_causal = causal;
|
||||
}
|
||||
|
||||
private:
|
||||
bool m_causal = false;
|
||||
};
|
||||
|
|
|
|||
|
|
@ -21,7 +21,7 @@ using namespace ov::op;
|
|||
|
||||
bool DictParameterResolver::run_on_model(const std::shared_ptr<Model>& model) {
|
||||
bool changed = false;
|
||||
const auto& parameters = model->get_parameters();
|
||||
const auto parameters = model->get_parameters();
|
||||
ParameterVector new_params;
|
||||
|
||||
for (const auto& p : parameters) {
|
||||
|
|
|
|||
|
|
@ -14,13 +14,12 @@
|
|||
OPENVINO_ASSERT(m_ptr != nullptr, "OpenVINO Runtime Plugin was not initialized."); \
|
||||
try { \
|
||||
__VA_ARGS__; \
|
||||
} catch (const ov::NotImplemented&) { \
|
||||
OPENVINO_NOT_IMPLEMENTED; \
|
||||
} catch (const std::exception& ex) { \
|
||||
OPENVINO_THROW(ex.what()); \
|
||||
} catch (...) { \
|
||||
OPENVINO_THROW("Unexpected exception"); \
|
||||
}
|
||||
|
||||
ov::Plugin::~Plugin() {
|
||||
m_ptr = {};
|
||||
}
|
||||
|
|
@ -46,7 +45,7 @@ const ov::Version ov::Plugin::get_version() const {
|
|||
}
|
||||
|
||||
void ov::Plugin::set_property(const ov::AnyMap& config) {
|
||||
OV_PLUGIN_CALL_STATEMENT(m_ptr->set_property(config));
|
||||
m_ptr->set_property(config);
|
||||
}
|
||||
|
||||
ov::SoPtr<ov::ICompiledModel> ov::Plugin::compile_model(const std::shared_ptr<const ov::Model>& model,
|
||||
|
|
@ -99,7 +98,7 @@ ov::SoPtr<ov::IRemoteContext> ov::Plugin::get_default_context(const AnyMap& para
|
|||
}
|
||||
|
||||
ov::Any ov::Plugin::get_property(const std::string& name, const AnyMap& arguments) const {
|
||||
OV_PLUGIN_CALL_STATEMENT({ return {m_ptr->get_property(name, arguments), {m_so}}; });
|
||||
return {m_ptr->get_property(name, arguments), {m_so}};
|
||||
}
|
||||
|
||||
bool ov::Plugin::supports_model_caching() const {
|
||||
|
|
|
|||
|
|
@ -378,6 +378,64 @@ TEST_F(TransformationTestsF, SmartReshapeReshapeAMatMulSeveralConsumers) {
|
|||
manager.register_pass<ov::pass::ReshapeAMatMul>();
|
||||
}
|
||||
|
||||
TEST_F(TransformationTestsF, SmartReshapeReshapeA_1DOtherInput) {
|
||||
{
|
||||
auto input_to_reshape = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::Shape{3, 2, 3});
|
||||
auto reshape_const = ov::op::v0::Constant::create(ov::element::i32, {2}, {3, 6});
|
||||
auto reshape = std::make_shared<ov::op::v1::Reshape>(input_to_reshape, reshape_const, false);
|
||||
|
||||
auto other_input = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::Shape{6});
|
||||
auto matmul = std::make_shared<ov::op::v0::MatMul>(reshape, other_input);
|
||||
model = std::make_shared<ov::Model>(ov::NodeVector{matmul}, ov::ParameterVector{input_to_reshape, other_input});
|
||||
manager.register_pass<ov::pass::ReshapeAMatMul>();
|
||||
}
|
||||
{
|
||||
auto input_to_reshape = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::Shape{3, 2, 3});
|
||||
auto other_input = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::Shape{6});
|
||||
const auto in_C_0 = std::make_shared<ov::op::v3::ShapeOf>(other_input);
|
||||
const auto in_C_1 = ov::op::v0::Constant::create(ov::element::i64, {1}, {-1});
|
||||
const auto in_C_2 = ov::op::v0::Constant::create(ov::element::i64, {}, {0});
|
||||
const auto C = std::make_shared<ov::op::v8::Gather>(in_C_0, in_C_1, in_C_2);
|
||||
|
||||
const auto N = ov::op::v0::Constant::create(ov::element::i64, {1}, {-1});
|
||||
const auto new_reshape_pattern = std::make_shared<ov::op::v0::Concat>(ov::OutputVector{N, C}, 0);
|
||||
auto reshape = std::make_shared<ov::op::v1::Reshape>(input_to_reshape, new_reshape_pattern, false);
|
||||
|
||||
auto matmul = std::make_shared<ov::op::v0::MatMul>(reshape, other_input);
|
||||
model_ref =
|
||||
std::make_shared<ov::Model>(ov::NodeVector{matmul}, ov::ParameterVector{input_to_reshape, other_input});
|
||||
}
|
||||
}
|
||||
|
||||
TEST_F(TransformationTestsF, SmartReshapeReshapeB_1DOtherInput) {
|
||||
{
|
||||
auto input_to_reshape = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::Shape{3, 2, 3});
|
||||
auto reshape_const = ov::op::v0::Constant::create(ov::element::i32, {2}, {3, 6});
|
||||
auto reshape = std::make_shared<ov::op::v1::Reshape>(input_to_reshape, reshape_const, false);
|
||||
|
||||
auto other_input = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::Shape{3});
|
||||
auto matmul = std::make_shared<ov::op::v0::MatMul>(other_input, reshape);
|
||||
model = std::make_shared<ov::Model>(ov::NodeVector{matmul}, ov::ParameterVector{input_to_reshape, other_input});
|
||||
manager.register_pass<ov::pass::ReshapeBMatMul>();
|
||||
}
|
||||
{
|
||||
auto input_to_reshape = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::Shape{3, 2, 3});
|
||||
auto other_input = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::Shape{3});
|
||||
const auto in_C_0 = std::make_shared<ov::op::v3::ShapeOf>(other_input);
|
||||
const auto in_C_1 = ov::op::v0::Constant::create(ov::element::i64, {1}, {-1});
|
||||
const auto in_C_2 = ov::op::v0::Constant::create(ov::element::i64, {}, {0});
|
||||
const auto C = std::make_shared<ov::op::v8::Gather>(in_C_0, in_C_1, in_C_2);
|
||||
|
||||
const auto N = ov::op::v0::Constant::create(ov::element::i64, {1}, {-1});
|
||||
const auto new_reshape_pattern = std::make_shared<ov::op::v0::Concat>(ov::OutputVector{C, N}, 0);
|
||||
auto reshape = std::make_shared<ov::op::v1::Reshape>(input_to_reshape, new_reshape_pattern, false);
|
||||
|
||||
auto matmul = std::make_shared<ov::op::v0::MatMul>(other_input, reshape);
|
||||
model_ref =
|
||||
std::make_shared<ov::Model>(ov::NodeVector{matmul}, ov::ParameterVector{input_to_reshape, other_input});
|
||||
}
|
||||
}
|
||||
|
||||
TEST_F(TransformationTestsF, SmartReshapeReshapeBMatMulSeveralConsumers) {
|
||||
// Reshape has 2 consumers: matmul and reduce.
|
||||
// Since reshape movement leads to loop creation (circular dependencies), the transformation can't be applied
|
||||
|
|
|
|||
|
|
@ -479,6 +479,140 @@ std::set<std::vector<element::Type>> jit_exp_emitter::get_supported_precisions(c
|
|||
return {{element::f32}};
|
||||
}
|
||||
|
||||
/// GELU_ERF ///
|
||||
jit_gelu_erf_emitter::jit_gelu_erf_emitter(dnnl::impl::cpu::aarch64::jit_generator* host,
|
||||
dnnl::impl::cpu::aarch64::cpu_isa_t host_isa,
|
||||
const std::shared_ptr<ov::Node>& node)
|
||||
: jit_emitter(host, host_isa, node, get_arithmetic_binary_exec_precision(node)) {
|
||||
prepare_table();
|
||||
exp_emitter = std::make_unique<jit_exp_emitter>(h, host_isa, node);
|
||||
}
|
||||
|
||||
jit_gelu_erf_emitter::jit_gelu_erf_emitter(dnnl::impl::cpu::aarch64::jit_generator* host,
|
||||
dnnl::impl::cpu::aarch64::cpu_isa_t host_isa,
|
||||
const ov::element::Type exec_prc) : jit_emitter(host, host_isa, exec_prc) {
|
||||
prepare_table();
|
||||
exp_emitter = std::make_unique<jit_exp_emitter>(h, host_isa, exec_prc);
|
||||
}
|
||||
|
||||
size_t jit_gelu_erf_emitter::get_inputs_count() const { return 1; }
|
||||
|
||||
size_t jit_gelu_erf_emitter::get_aux_vecs_count() const {
|
||||
return std::max<size_t>(exp_emitter->get_aux_vecs_count() + 3, 7);
|
||||
}
|
||||
|
||||
size_t jit_gelu_erf_emitter::get_aux_gprs_count() const {
|
||||
return exp_emitter->get_aux_gprs_count() + 1;
|
||||
}
|
||||
|
||||
void jit_gelu_erf_emitter::emit_impl(const std::vector<size_t> &in_vec_idxs, const std::vector<size_t> &out_vec_idxs) const {
|
||||
if (host_isa_ == dnnl::impl::cpu::aarch64::asimd) {
|
||||
emit_isa<dnnl::impl::cpu::aarch64::asimd>(in_vec_idxs, out_vec_idxs);
|
||||
} else {
|
||||
OV_CPU_JIT_EMITTER_THROW("Can't create jit eltwise kernel");
|
||||
}
|
||||
}
|
||||
|
||||
template <dnnl::impl::cpu::aarch64::cpu_isa_t isa>
|
||||
void jit_gelu_erf_emitter::emit_isa(const std::vector<size_t> &in_vec_idxs, const std::vector<size_t> &out_vec_idxs) const {
|
||||
OV_CPU_JIT_EMITTER_ASSERT(exec_prc_ == ov::element::f32, "unsupported precision: " + exec_prc_.to_string());
|
||||
|
||||
using TReg = typename dnnl::impl::cpu::aarch64::cpu_isa_traits<isa>::TReg;
|
||||
const TReg vmm_src(in_vec_idxs[0]);
|
||||
const TReg vmm_dst(out_vec_idxs[0]);
|
||||
|
||||
const TReg vmm_aux0(aux_vec_idxs[0]);
|
||||
const TReg vmm_aux1(aux_vec_idxs[1]);
|
||||
const TReg vmm_aux2(aux_vec_idxs[2]);
|
||||
const TReg vmm_aux3(aux_vec_idxs[3]);
|
||||
const TReg vmm_aux(aux_vec_idxs[std::max<size_t>(exp_emitter->get_aux_vecs_count(), 4)]);
|
||||
const TReg vmm_aux_t(aux_vec_idxs[std::max<size_t>(exp_emitter->get_aux_vecs_count() + 1, 5)]);
|
||||
const TReg vmm_aux_dst(aux_vec_idxs[std::max<size_t>(exp_emitter->get_aux_vecs_count() + 2, 6)]);
|
||||
|
||||
// x = s / sqrt(2)
|
||||
h->ld1r(vmm_aux0.s, table_val2("gelu_erf_one_over_sqrt_two"));
|
||||
h->fmul(vmm_aux0.s, vmm_aux0.s, vmm_src.s);
|
||||
|
||||
// abs(x)
|
||||
h->fabs(vmm_aux0.s, vmm_aux0.s);
|
||||
|
||||
// t = 1 / (p*x + 1)
|
||||
h->ld1r(vmm_aux1.s, table_val2("gelu_erf_approx_const"));
|
||||
h->ld1r(vmm_aux2.s, table_val2("one"));
|
||||
h->mov(vmm_aux3.b16, vmm_aux2.b16);
|
||||
h->fmla(vmm_aux2.s, vmm_aux1.s, vmm_aux0.s);
|
||||
h->fdiv(vmm_aux_t.s, vmm_aux3.s, vmm_aux2.s);
|
||||
|
||||
// -exp(-x*x)
|
||||
h->fmul(vmm_aux.s, vmm_aux0.s, vmm_aux0.s);
|
||||
h->ld1r(vmm_aux2.s, table_val2("sign_mask"));
|
||||
h->orr(vmm_aux.b16, vmm_aux.b16, vmm_aux2.b16);
|
||||
exp_emitter->emit_code(
|
||||
{ vmm_aux.getIdx() },
|
||||
{ vmm_aux_dst.getIdx() },
|
||||
aux_vec_idxs,
|
||||
aux_gpr_idxs);
|
||||
h->ld1r(vmm_aux2.s, table_val2("sign_mask"));
|
||||
// vmm_aux_dst = -exp(-x*x)
|
||||
h->orr(vmm_aux_dst.b16, vmm_aux_dst.b16, vmm_aux2.b16);
|
||||
|
||||
// get sign
|
||||
h->and_(vmm_aux.b16, vmm_src.b16, vmm_aux2.b16);
|
||||
|
||||
// -exp(-x*x)*t
|
||||
h->fmul(vmm_aux_dst.s, vmm_aux_dst.s, vmm_aux_t.s);
|
||||
|
||||
// compute polynomialial r
|
||||
h->ld1r(vmm_aux0.s, table_val2("erf_pol5"));
|
||||
h->ld1r(vmm_aux1.s, table_val2("erf_pol4"));
|
||||
h->fmla(vmm_aux1.s, vmm_aux0.s, vmm_aux_t.s);
|
||||
|
||||
h->ld1r(vmm_aux0.s, table_val2("erf_pol3"));
|
||||
h->fmla(vmm_aux0.s, vmm_aux1.s, vmm_aux_t.s);
|
||||
|
||||
h->ld1r(vmm_aux1.s, table_val2("erf_pol2"));
|
||||
h->fmla(vmm_aux1.s, vmm_aux0.s, vmm_aux_t.s);
|
||||
|
||||
h->ld1r(vmm_aux0.s, table_val2("erf_pol1"));
|
||||
h->fmla(vmm_aux0.s, vmm_aux1.s, vmm_aux_t.s);
|
||||
|
||||
// erf = sign * (1 - r * t * exp(-x*x))
|
||||
h->ld1r(vmm_aux2.s, table_val2("one"));
|
||||
h->fmla(vmm_aux2.s, vmm_aux0.s, vmm_aux_dst.s);
|
||||
h->orr(vmm_aux2.b16, vmm_aux.b16, vmm_aux2.b16);
|
||||
|
||||
// S = 0.5 * s
|
||||
h->ld1r(vmm_aux3.s, table_val2("half"));
|
||||
h->fmul(vmm_dst.s, vmm_src.s, vmm_aux3.s);
|
||||
// GELU = 0.5 * s * (1 + erf) = S + S * erf
|
||||
h->fmla(vmm_dst.s, vmm_dst.s, vmm_aux2.s);
|
||||
}
|
||||
|
||||
void jit_gelu_erf_emitter::register_table_entries() {
|
||||
push_arg_entry_of("one", 0x3f800000, true);
|
||||
push_arg_entry_of("half", 0x3f000000, true);
|
||||
push_arg_entry_of("sign_mask", 0x80000000, true);
|
||||
|
||||
push_arg_entry_of("gelu_erf_approx_const", 0x3ea7ba05, true);
|
||||
push_arg_entry_of("gelu_erf_one_over_sqrt_two", 0x3f3504f3, true);
|
||||
push_arg_entry_of("gelu_erf_one_over_sqrt_pi", 0x3f106eba, true);
|
||||
|
||||
push_arg_entry_of("erf_pol1", 0x3e827906, true); // p1 = 0.254829592f
|
||||
push_arg_entry_of("erf_pol2", 0xbe91a98e, true); // p2 = -0.284496736f
|
||||
push_arg_entry_of("erf_pol3", 0x3fb5f0e3, true); // p3 = 1.421413741f
|
||||
push_arg_entry_of("erf_pol4", 0xbfba00e3, true); // p4 = -1.453152027f
|
||||
push_arg_entry_of("erf_pol5", 0x3f87dc22, true); // p5 = 1.061405429f
|
||||
}
|
||||
|
||||
void jit_gelu_erf_emitter::emit_data() const {
|
||||
jit_emitter::emit_data();
|
||||
exp_emitter->emit_data();
|
||||
}
|
||||
|
||||
std::set<std::vector<element::Type>> jit_gelu_erf_emitter::get_supported_precisions(const std::shared_ptr<ov::Node>& node) {
|
||||
return {{element::f32}};
|
||||
}
|
||||
|
||||
/// HARD_SWISH ///
|
||||
jit_hswish_emitter::jit_hswish_emitter(dnnl::impl::cpu::aarch64::jit_generator* host,
|
||||
dnnl::impl::cpu::aarch64::cpu_isa_t host_isa,
|
||||
|
|
|
|||
|
|
@ -193,6 +193,37 @@ private:
|
|||
void emit_isa(const std::vector<size_t> &in_vec_idxs, const std::vector<size_t> &out_vec_idxs) const;
|
||||
};
|
||||
|
||||
class jit_gelu_erf_emitter : public jit_emitter {
|
||||
public:
|
||||
jit_gelu_erf_emitter(dnnl::impl::cpu::aarch64::jit_generator* host,
|
||||
dnnl::impl::cpu::aarch64::cpu_isa_t host_isa,
|
||||
const ov::element::Type exec_prc = ov::element::f32);
|
||||
|
||||
jit_gelu_erf_emitter(dnnl::impl::cpu::aarch64::jit_generator* host,
|
||||
dnnl::impl::cpu::aarch64::cpu_isa_t host_isa,
|
||||
const std::shared_ptr<ov::Node>& node);
|
||||
|
||||
size_t get_inputs_count() const override;
|
||||
|
||||
size_t get_aux_vecs_count() const override;
|
||||
|
||||
size_t get_aux_gprs_count() const override;
|
||||
|
||||
void register_table_entries() override;
|
||||
|
||||
void emit_data() const override;
|
||||
|
||||
static std::set<std::vector<element::Type>> get_supported_precisions(const std::shared_ptr<ov::Node>& node = nullptr);
|
||||
|
||||
private:
|
||||
std::unique_ptr<jit_exp_emitter> exp_emitter;
|
||||
|
||||
void emit_impl(const std::vector<size_t> &in_vec_idxs, const std::vector<size_t> &out_vec_idxs) const override;
|
||||
|
||||
template <dnnl::impl::cpu::aarch64::cpu_isa_t isa>
|
||||
void emit_isa(const std::vector<size_t> &in_vec_idxs, const std::vector<size_t> &out_vec_idxs) const;
|
||||
};
|
||||
|
||||
class jit_hswish_emitter : public jit_emitter {
|
||||
public:
|
||||
jit_hswish_emitter(dnnl::impl::cpu::aarch64::jit_generator *host,
|
||||
|
|
|
|||
|
|
@ -382,8 +382,7 @@ struct ConvertPrecision<std::tuple<src_t, ov::float16>> {
|
|||
src_t lbound, ubound;
|
||||
std::tie(lbound, ubound) = ctx.range<src_t>();
|
||||
|
||||
if (std::is_integral<src_t>::value
|
||||
|| ctx.interimPrc.is_real()) {
|
||||
if (std::is_integral<src_t>::value) {
|
||||
parallel_for(iterations, [&](size_t i) {
|
||||
batch_type tmp;
|
||||
const size_t offset = i * batch;
|
||||
|
|
@ -392,6 +391,19 @@ struct ConvertPrecision<std::tuple<src_t, ov::float16>> {
|
|||
tmp[j] = static_cast<float>(std::max(std::min(src[offset + j], ubound), lbound));
|
||||
jit_convert(tmp, dst + offset, current_batch_size); // fp32 -> fp16
|
||||
});
|
||||
} else if (ctx.interimPrc.is_real()) {
|
||||
parallel_for(iterations, [&](size_t i) {
|
||||
const size_t offset = i * batch;
|
||||
const size_t current_batch_size = std::min(ctx.size - offset, batch);
|
||||
if (std::is_same<typename std::remove_cv<src_t>::type, float>::value) { // fp32 -> fp16
|
||||
jit_convert(reinterpret_cast<const float *>(src) + offset, dst + offset, current_batch_size);
|
||||
} else {
|
||||
batch_type tmp;
|
||||
for (size_t j = 0; j < current_batch_size; ++j) // src_t -> fp32
|
||||
tmp[j] = static_cast<float>(src[offset + j]);
|
||||
jit_convert(tmp, dst + offset, current_batch_size); // fp32 -> fp16
|
||||
}
|
||||
});
|
||||
} else {
|
||||
parallel_for(iterations, [&](size_t i) {
|
||||
batch_type tmp;
|
||||
|
|
@ -420,8 +432,7 @@ struct ConvertPrecision<std::tuple<ov::float16, dst_t>> {
|
|||
float lbound, ubound;
|
||||
std::tie(lbound, ubound) = ctx.range<ov::float16>();
|
||||
|
||||
if (ctx.interimPrc.is_real()
|
||||
|| std::is_integral<dst_t>::value) {
|
||||
if (std::is_integral<dst_t>::value) {
|
||||
parallel_for(iterations, [&](size_t i) {
|
||||
batch_type tmp;
|
||||
const size_t offset = i * batch;
|
||||
|
|
@ -430,6 +441,19 @@ struct ConvertPrecision<std::tuple<ov::float16, dst_t>> {
|
|||
for (size_t j = 0; j < current_batch_size; ++j) // fp32 -> dst_t
|
||||
dst[offset + j] = static_cast<dst_t>(std::max(std::min(tmp[j], ubound), lbound));
|
||||
});
|
||||
} else if (ctx.interimPrc.is_real()) {
|
||||
parallel_for(iterations, [&](size_t i) {
|
||||
const size_t offset = i * batch;
|
||||
const size_t current_batch_size = std::min(ctx.size - offset, batch);
|
||||
if (std::is_same<typename std::remove_cv<dst_t>::type, float>::value) { // fp16 -> fp32
|
||||
jit_convert(src + offset, reinterpret_cast<float *>(dst) + offset, current_batch_size);
|
||||
} else {
|
||||
batch_type tmp;
|
||||
jit_convert(src + offset, tmp, current_batch_size); // fp16 -> fp32
|
||||
for (size_t j = 0; j < current_batch_size; ++j) // fp32 -> dst_t
|
||||
dst[offset + j] = static_cast<dst_t>(tmp[j]);
|
||||
}
|
||||
});
|
||||
} else {
|
||||
parallel_for(iterations, [&](size_t i) {
|
||||
batch_type tmp;
|
||||
|
|
|
|||
|
|
@ -25,6 +25,7 @@ bool JitEltwiseExecutor::isSupported(
|
|||
Algorithm::EltwiseElu,
|
||||
Algorithm::EltwiseEqual,
|
||||
Algorithm::EltwiseExp,
|
||||
Algorithm::EltwiseGeluErf,
|
||||
Algorithm::EltwiseHswish,
|
||||
Algorithm::EltwiseMaximum,
|
||||
Algorithm::EltwiseMinimum,
|
||||
|
|
|
|||
|
|
@ -640,6 +640,7 @@ std::shared_ptr<jit_emitter> jit_uni_eltwise_generic<isa>::create_eltwise_emitte
|
|||
OV_CASE(Algorithm::EltwiseMaximum, ov::intel_cpu::aarch64::jit_maximum_emitter),
|
||||
OV_CASE(Algorithm::EltwiseMinimum, ov::intel_cpu::aarch64::jit_minimum_emitter),
|
||||
OV_CASE(Algorithm::EltwiseMish, ov::intel_cpu::aarch64::jit_mish_emitter),
|
||||
OV_CASE(Algorithm::EltwiseGeluErf, ov::intel_cpu::aarch64::jit_gelu_erf_emitter),
|
||||
OV_CASE(Algorithm::EltwiseMulAdd, ov::intel_cpu::aarch64::jit_mul_add_emitter),
|
||||
OV_CASE(Algorithm::EltwiseMod, ov::intel_cpu::aarch64::jit_mod_emitter),
|
||||
OV_CASE(Algorithm::EltwiseMultiply, ov::intel_cpu::aarch64::jit_multiply_emitter),
|
||||
|
|
@ -803,6 +804,7 @@ std::set<std::vector<element::Type>> eltwise_precision_helper::get_supported_pre
|
|||
OV_CASE(Algorithm::EltwiseElu, jit_elu_emitter),
|
||||
OV_CASE(Algorithm::EltwiseEqual, jit_equal_emitter),
|
||||
OV_CASE(Algorithm::EltwiseExp, jit_exp_emitter),
|
||||
OV_CASE(Algorithm::EltwiseGeluErf, jit_gelu_erf_emitter),
|
||||
OV_CASE(Algorithm::EltwiseHswish, jit_hswish_emitter),
|
||||
OV_CASE(Algorithm::EltwiseMaximum, jit_maximum_emitter),
|
||||
OV_CASE(Algorithm::EltwiseMinimum, jit_minimum_emitter),
|
||||
|
|
|
|||
|
|
@ -160,6 +160,7 @@ std::string ActivationLayerCPUTest::getPrimitiveType(const utils::ActivationType
|
|||
(activation_type == utils::ActivationTypes::HSwish) ||
|
||||
(activation_type == utils::ActivationTypes::HardSigmoid) ||
|
||||
(activation_type == utils::ActivationTypes::Mish) ||
|
||||
(activation_type == utils::ActivationTypes::GeluErf) ||
|
||||
(activation_type == utils::ActivationTypes::Relu) ||
|
||||
(activation_type == utils::ActivationTypes::Sigmoid) ||
|
||||
(activation_type == utils::ActivationTypes::Swish) ||
|
||||
|
|
|
|||
|
|
@ -0,0 +1,78 @@
|
|||
// Copyright (C) 2024 Intel Corporation
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
//
|
||||
|
||||
#pragma once
|
||||
|
||||
#include "intel_gpu/op/sdpa.hpp"
|
||||
#include "openvino/core/node.hpp"
|
||||
#include "openvino/core/partial_shape.hpp"
|
||||
#include "openvino/op/op.hpp"
|
||||
|
||||
namespace ov {
|
||||
namespace intel_gpu {
|
||||
namespace op {
|
||||
|
||||
class IndirectSDPA : public ov::intel_gpu::op::SDPA {
|
||||
public:
|
||||
OPENVINO_OP("IndirectSDPA", "gpu_opset");
|
||||
|
||||
IndirectSDPA() = default;
|
||||
|
||||
IndirectSDPA(const ov::Output<Node>& Q,
|
||||
const ov::Output<Node>& K,
|
||||
const ov::Output<Node>& V,
|
||||
const ov::Output<Node>& beam_table,
|
||||
const bool is_causal,
|
||||
const int64_t indirect_axis,
|
||||
const std::vector<int64_t>& order_q,
|
||||
const std::vector<int64_t>& order_k,
|
||||
const std::vector<int64_t>& order_v,
|
||||
const std::vector<int64_t>& order_out,
|
||||
const ov::element::Type output_type = ov::element::undefined);
|
||||
|
||||
IndirectSDPA(const ov::Output<Node>& Q,
|
||||
const ov::Output<Node>& K,
|
||||
const ov::Output<Node>& V,
|
||||
const ov::Output<Node>& attn_mask,
|
||||
const ov::Output<Node>& beam_table,
|
||||
const bool is_causal,
|
||||
const int64_t indirect_axis,
|
||||
const std::vector<int64_t>& order_q,
|
||||
const std::vector<int64_t>& order_k,
|
||||
const std::vector<int64_t>& order_v,
|
||||
const std::vector<int64_t>& order_out,
|
||||
const ov::element::Type output_type = ov::element::undefined);
|
||||
|
||||
IndirectSDPA(const ov::Output<Node>& Q,
|
||||
const ov::Output<Node>& K,
|
||||
const ov::Output<Node>& V,
|
||||
const ov::Output<Node>& attn_mask,
|
||||
const ov::Output<Node>& scale,
|
||||
const ov::Output<Node>& beam_table,
|
||||
const bool is_causal,
|
||||
const int64_t indirect_axis,
|
||||
const std::vector<int64_t>& order_q,
|
||||
const std::vector<int64_t>& order_k,
|
||||
const std::vector<int64_t>& order_v,
|
||||
const std::vector<int64_t>& order_out,
|
||||
const ov::element::Type output_type = ov::element::undefined);
|
||||
|
||||
bool visit_attributes(ov::AttributeVisitor &visitor) override;
|
||||
void validate_and_infer_types() override;
|
||||
|
||||
std::shared_ptr<Node> clone_with_new_inputs(const ov::OutputVector& new_args) const override;
|
||||
|
||||
ov::element::Type get_output_type() const { return m_output_type; }
|
||||
|
||||
int64_t get_indirect_axis() const { return m_indirect_axis; }
|
||||
|
||||
using ov::intel_gpu::op::SDPA::default_order;
|
||||
|
||||
protected:
|
||||
int64_t m_indirect_axis = -1;
|
||||
};
|
||||
|
||||
} // namespace op
|
||||
} // namespace intel_gpu
|
||||
} // namespace ov
|
||||
|
|
@ -285,3 +285,4 @@ REGISTER_FACTORY(internal, IndirectGemm);
|
|||
REGISTER_FACTORY(internal, Convolution);
|
||||
REGISTER_FACTORY(internal, Placeholder);
|
||||
REGISTER_FACTORY(internal, SDPA);
|
||||
REGISTER_FACTORY(internal, IndirectSDPA);
|
||||
|
|
|
|||
|
|
@ -19,6 +19,7 @@ struct scaled_dot_product_attention : public primitive_base<scaled_dot_product_a
|
|||
scaled_dot_product_attention(const primitive_id& id,
|
||||
const std::vector<cldnn::input_info> inputs,
|
||||
bool is_causal,
|
||||
int64_t indirect_axis = -1,
|
||||
const std::vector<int64_t>& input_q_transpose_order = {},
|
||||
const std::vector<int64_t>& input_k_transpose_order = {},
|
||||
const std::vector<int64_t>& input_v_transpose_order = {},
|
||||
|
|
@ -26,17 +27,23 @@ struct scaled_dot_product_attention : public primitive_base<scaled_dot_product_a
|
|||
const padding& output_padding = padding())
|
||||
: primitive_base(id, inputs, {output_padding})
|
||||
, is_causal(is_causal)
|
||||
, has_attn_mask_input(inputs.size() > 3)
|
||||
, has_scale_input(inputs.size() > 4)
|
||||
, indirect_axis(indirect_axis)
|
||||
, input_q_transpose_order(input_q_transpose_order)
|
||||
, input_k_transpose_order(input_k_transpose_order)
|
||||
, input_v_transpose_order(input_v_transpose_order)
|
||||
, output_transpose_order(output_transpose_order) {}
|
||||
, output_transpose_order(output_transpose_order) {
|
||||
auto data_inputs_num = inputs.size();
|
||||
if (indirect_axis != -1)
|
||||
data_inputs_num--;
|
||||
|
||||
has_attn_mask_input = data_inputs_num > 3;
|
||||
has_scale_input = data_inputs_num > 4;
|
||||
}
|
||||
|
||||
bool is_causal = false;
|
||||
bool has_attn_mask_input = false;
|
||||
bool has_scale_input = false;
|
||||
int64_t indirect_axis = -1;
|
||||
|
||||
std::vector<int64_t> input_q_transpose_order;
|
||||
std::vector<int64_t> input_k_transpose_order;
|
||||
|
|
@ -48,6 +55,7 @@ struct scaled_dot_product_attention : public primitive_base<scaled_dot_product_a
|
|||
seed = hash_combine(seed, is_causal);
|
||||
seed = hash_combine(seed, has_attn_mask_input);
|
||||
seed = hash_combine(seed, has_scale_input);
|
||||
seed = hash_combine(seed, indirect_axis);
|
||||
seed = hash_range(seed, input_q_transpose_order.begin(), input_q_transpose_order.end());
|
||||
seed = hash_range(seed, input_k_transpose_order.begin(), input_k_transpose_order.end());
|
||||
seed = hash_range(seed, input_v_transpose_order.begin(), input_v_transpose_order.end());
|
||||
|
|
@ -64,6 +72,7 @@ struct scaled_dot_product_attention : public primitive_base<scaled_dot_product_a
|
|||
return is_causal == rhs_casted.is_causal &&
|
||||
has_attn_mask_input == rhs_casted.has_attn_mask_input &&
|
||||
has_scale_input == rhs_casted.has_scale_input &&
|
||||
indirect_axis == rhs_casted.indirect_axis &&
|
||||
input_q_transpose_order == rhs_casted.input_q_transpose_order &&
|
||||
input_k_transpose_order == rhs_casted.input_k_transpose_order &&
|
||||
input_v_transpose_order == rhs_casted.input_v_transpose_order &&
|
||||
|
|
@ -75,6 +84,7 @@ struct scaled_dot_product_attention : public primitive_base<scaled_dot_product_a
|
|||
ob << is_causal;
|
||||
ob << has_attn_mask_input;
|
||||
ob << has_scale_input;
|
||||
ob << indirect_axis;
|
||||
ob << input_q_transpose_order;
|
||||
ob << input_k_transpose_order;
|
||||
ob << input_v_transpose_order;
|
||||
|
|
@ -86,6 +96,7 @@ struct scaled_dot_product_attention : public primitive_base<scaled_dot_product_a
|
|||
ib >> is_causal;
|
||||
ib >> has_attn_mask_input;
|
||||
ib >> has_scale_input;
|
||||
ib >> indirect_axis;
|
||||
ib >> input_q_transpose_order;
|
||||
ib >> input_k_transpose_order;
|
||||
ib >> input_v_transpose_order;
|
||||
|
|
|
|||
|
|
@ -129,6 +129,7 @@ public:
|
|||
std::vector<std::string> forced_impl_types; // Force implementation type either ocl or onednn
|
||||
int max_kernels_per_batch; // Maximum number of kernels in a batch during compiling kernels
|
||||
int impls_cache_capacity; // The maximum number of entries in the kernel impl cache
|
||||
int enable_sdpa; // Allows to control SDPA decomposition
|
||||
int disable_async_compilation; // Disable async compilation
|
||||
int disable_winograd_conv; // Disable Winograd conv
|
||||
int disable_dynamic_impl; // Disable dynamic implementation
|
||||
|
|
|
|||
|
|
@ -2,34 +2,188 @@
|
|||
// SPDX-License-Identifier: Apache-2.0
|
||||
//
|
||||
|
||||
#include "primitive_base.hpp"
|
||||
#include "multi_stage_primitive.hpp"
|
||||
#include "scaled_dot_product_attention_inst.h"
|
||||
#include "kv_cache_inst.h"
|
||||
|
||||
#include "sdpa/sdpa_kernel_selector.h"
|
||||
#include "sdpa/sdpa_kernel_base.h"
|
||||
|
||||
namespace cldnn {
|
||||
namespace ocl {
|
||||
struct scaled_dot_product_attention_impl : typed_primitive_impl_ocl<scaled_dot_product_attention> {
|
||||
using parent = typed_primitive_impl_ocl<scaled_dot_product_attention>;
|
||||
|
||||
// SDPA impl may create 2 versions of the kernel internally
|
||||
// 1. Default SDPA kernels
|
||||
// 2. SDPA kernels with indirect access to one of the inputs
|
||||
// This feature is used to avoid perf drop when we create single kernel which checks batch size in runtime
|
||||
// Can be reverted once performance of the kernel is improved
|
||||
struct scaled_dot_product_attention_impl : multi_stage_primitive<scaled_dot_product_attention> {
|
||||
using parent = multi_stage_primitive<scaled_dot_product_attention>;
|
||||
using parent::parent;
|
||||
using kernel_selector_t = kernel_selector::sdpa_kernel_selector;
|
||||
using kernel_params_t = kernel_selector::sdpa_params;
|
||||
|
||||
DECLARE_OBJECT_TYPE_SERIALIZATION(cldnn::ocl::scaled_dot_product_attention_impl)
|
||||
|
||||
const uint32_t default_sdpa = 0;
|
||||
const uint32_t indirect_sdpa = 1;
|
||||
|
||||
std::unique_ptr<primitive_impl> clone() const override {
|
||||
return make_unique<scaled_dot_product_attention_impl>(*this);
|
||||
}
|
||||
|
||||
scaled_dot_product_attention_impl() = default;
|
||||
|
||||
scaled_dot_product_attention_impl(const std::vector<kernel_selector::kernel_data>& kd) : parent(kd) {
|
||||
this->can_reuse_memory = true;
|
||||
}
|
||||
|
||||
void load(BinaryInputBuffer& ib) override {
|
||||
parent::load(ib);
|
||||
if (is_dynamic()) {
|
||||
auto& kernel_selector = kernel_selector_t::Instance();
|
||||
auto kernel_impl = kernel_selector.GetImplementation(_kernel_data.kernelName);
|
||||
kernel_impl->GetUpdateDispatchDataFunc(_kernel_data);
|
||||
auto kernel_impl = kernel_selector.GetImplementation(_kernels_data[default_sdpa].kernelName);
|
||||
kernel_impl->GetUpdateDispatchDataFunc(_kernels_data[default_sdpa]);
|
||||
if (_kernels_data.size() == 2) {
|
||||
auto bt_kernel_impl = kernel_selector.GetImplementation(_kernels_data[indirect_sdpa].kernelName);
|
||||
bt_kernel_impl->GetUpdateDispatchDataFunc(_kernels_data[indirect_sdpa]);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
protected:
|
||||
std::vector<layout> get_internal_buffer_layouts_impl() const override {
|
||||
// TODO: current implementation is supposed to have the same kernel version for both indirect/default paths,
|
||||
// considering this, we may assume that both indirect/default kernels have absolutely the same intermediate
|
||||
// buffers number and its' sizes (since update_dispatch_data is called for both kernels too), and
|
||||
// do not double memory allocations during reallocate_if_needed() function call
|
||||
std::vector<layout> layouts;
|
||||
if (_kernels_data.size() > 0) {
|
||||
auto dtype = from_data_type(_kernels_data[0].internalBufferDataType);
|
||||
const auto bpp = data_type_traits::size_of(dtype);
|
||||
for (auto size : _kernels_data[0].internalBufferSizes) {
|
||||
layout inbuf_layout = {dtype, format::bfyx, // simple linear format (flattern to x channel)
|
||||
{1, 1, 1, (tensor::value_type)(size / bpp)}};
|
||||
layouts.push_back(inbuf_layout);
|
||||
}
|
||||
}
|
||||
|
||||
return layouts;
|
||||
}
|
||||
|
||||
static size_t get_beam_table_id(std::shared_ptr<const scaled_dot_product_attention> primitive) {
|
||||
GPU_DEBUG_TRACE << "get_beam_table_id " << primitive->input_size() - 1 << "\n";
|
||||
return primitive->input_size() - 1;
|
||||
}
|
||||
|
||||
static bool has_indirect_inputs(const kernel_impl_params& impl_param) {
|
||||
const auto& desc = impl_param.typed_desc<scaled_dot_product_attention>();
|
||||
return desc->indirect_axis != -1;
|
||||
}
|
||||
|
||||
kernel_arguments_data get_arguments(const scaled_dot_product_attention_inst& instance, size_t stage) const override {
|
||||
kernel_arguments_data args;
|
||||
|
||||
auto inputs_num = instance.inputs_memory_count();
|
||||
if (instance.has_indirect_inputs() && stage == default_sdpa)
|
||||
inputs_num--;
|
||||
|
||||
for (size_t i = 0; i < inputs_num; i++) {
|
||||
args.inputs.push_back(instance.input_memory_ptr(i));
|
||||
}
|
||||
|
||||
if (instance.has_fused_primitives()) {
|
||||
size_t count = instance.get_fused_mem_count();
|
||||
for (size_t i = 0; i < count; i++) {
|
||||
args.fused_op_inputs.push_back(instance.fused_memory(i));
|
||||
}
|
||||
}
|
||||
|
||||
for (size_t i = 0; i < instance.outputs_memory_count(); i++) {
|
||||
args.outputs.push_back(instance.output_memory_ptr(i));
|
||||
}
|
||||
|
||||
args.shape_info = instance.shape_info_memory_ptr();
|
||||
|
||||
return args;
|
||||
}
|
||||
|
||||
void set_arguments_impl(scaled_dot_product_attention_inst& instance) override {}
|
||||
|
||||
event::ptr execute_stage(const std::vector<event::ptr>& events, scaled_dot_product_attention_inst& instance, size_t stage) {
|
||||
stream& stream = instance.get_network().get_stream();
|
||||
std::vector<event::ptr> tmp_events(events);
|
||||
std::vector<event::ptr> all_events;
|
||||
size_t kernel_offset = 0;
|
||||
|
||||
for (size_t s = 0; s < stage; s++) {
|
||||
kernel_offset += _kernels_data[s].kernels.size();
|
||||
}
|
||||
for (size_t kd_idx = 0; kd_idx < _kernels_data[stage].kernels.size(); ++kd_idx) {
|
||||
if (_kernels_data[stage].kernels[kd_idx].skip_execution)
|
||||
continue;
|
||||
|
||||
size_t idx_final = kernel_offset + kd_idx;
|
||||
// If any user of the desc's users is CPU implementation or network's output, set desc as a output event (event won't be nullptr)
|
||||
bool needs_completion_event = instance.needs_completion_event();
|
||||
|
||||
auto& params = _kernels_data[stage].kernels[kd_idx].params;
|
||||
auto args = get_arguments(instance, stage);
|
||||
args.scalars = ¶ms.scalars;
|
||||
|
||||
for (size_t i = 0; i < instance.get_intermediates_memories().size(); i++)
|
||||
args.intermediates.push_back(instance.get_intermediates_memories()[i]);
|
||||
|
||||
stream.set_arguments(*_kernels[idx_final], _kernels_data[stage].kernels[kd_idx].params, args);
|
||||
|
||||
const auto& gws = params.workGroups.global;
|
||||
const auto& lws = params.workGroups.local;
|
||||
|
||||
GPU_DEBUG_TRACE_DETAIL << "Enqueue stage " << stage << " kernel " << idx_final << ": gws=[" << gws[0] << ", " << gws[1] << ", " << gws[2] << "] "
|
||||
<< "lws=[" << lws[0] << ", " << lws[1] << ", " << lws[2] << "]"
|
||||
<< (needs_completion_event ? " has_completion_event=true" : "") << std::endl;
|
||||
|
||||
auto ev = stream.enqueue_kernel(*_kernels[idx_final], params, args, tmp_events, needs_completion_event);
|
||||
if (_kernels_data[stage].needs_sub_kernels_sync) {
|
||||
tmp_events = {ev};
|
||||
}
|
||||
all_events.push_back(ev);
|
||||
}
|
||||
|
||||
return aggregate_events(all_events, stream, all_events.size() > 1);
|
||||
}
|
||||
|
||||
bool need_indirect_load(const scaled_dot_product_attention_inst& instance) const {
|
||||
auto desc = instance.get_typed_desc<scaled_dot_product_attention>();
|
||||
|
||||
if (!instance.has_indirect_inputs())
|
||||
return false;
|
||||
|
||||
const auto& params = *instance.get_impl_params();
|
||||
const auto indirect_axis = desc->indirect_axis;
|
||||
if (params.input_layouts[get_beam_table_id(desc)].get_partial_shape()[indirect_axis].get_length() == 1)
|
||||
return false;
|
||||
|
||||
const auto& deps = instance.dependencies();
|
||||
|
||||
const auto indirect_dep_idx = 1;
|
||||
const auto& indirect_dep = deps[indirect_dep_idx].first;
|
||||
if (dynamic_cast<const kv_cache_inst*>(indirect_dep) == nullptr) {
|
||||
return true;
|
||||
}
|
||||
|
||||
auto state_layout = indirect_dep->get_impl_params()->get_input_layout(0);
|
||||
bool is_prefill = state_layout.count() == 0;
|
||||
return !is_prefill;
|
||||
}
|
||||
|
||||
event::ptr execute_impl(const std::vector<event::ptr>& events, scaled_dot_product_attention_inst& instance) override {
|
||||
if (need_indirect_load(instance))
|
||||
return execute_stage(events, instance, indirect_sdpa);
|
||||
else
|
||||
return execute_stage(events, instance, default_sdpa);
|
||||
}
|
||||
|
||||
static kernel_selector::sdpa_configuration get_sdpa_configuration(const kernel_impl_params& impl_param) {
|
||||
kernel_selector::sdpa_configuration config;
|
||||
|
||||
|
|
@ -44,16 +198,16 @@ struct scaled_dot_product_attention_impl : typed_primitive_impl_ocl<scaled_dot_p
|
|||
return transposed_pshape;
|
||||
};
|
||||
|
||||
const auto& prim = impl_param.typed_desc<scaled_dot_product_attention>();
|
||||
const auto query_shape = transpose_pshape(impl_param.get_input_layout(0).get_partial_shape(), prim->input_q_transpose_order);
|
||||
const auto key_shape = transpose_pshape(impl_param.get_input_layout(1).get_partial_shape(), prim->input_k_transpose_order);
|
||||
const auto value_shape = transpose_pshape(impl_param.get_input_layout(2).get_partial_shape(), prim->input_v_transpose_order);
|
||||
const auto& desc = impl_param.typed_desc<scaled_dot_product_attention>();
|
||||
const auto query_shape = transpose_pshape(impl_param.get_input_layout(0).get_partial_shape(), desc->input_q_transpose_order);
|
||||
const auto key_shape = transpose_pshape(impl_param.get_input_layout(1).get_partial_shape(), desc->input_k_transpose_order);
|
||||
const auto value_shape = transpose_pshape(impl_param.get_input_layout(2).get_partial_shape(), desc->input_v_transpose_order);
|
||||
|
||||
OPENVINO_ASSERT(key_shape == value_shape, "[GPU] The shapes of key and value inputs are expected to be equal");
|
||||
for (size_t i = 0; i < query_shape.size(); ++i) {
|
||||
if (query_shape[i].is_static() && key_shape[i].is_static() && value_shape[i].is_static()) {
|
||||
if (query_shape[i].get_length() > key_shape[i].get_length()) {
|
||||
config.broadcast_axis = prim->input_k_transpose_order[i];
|
||||
config.broadcast_axis = desc->input_k_transpose_order[i];
|
||||
config.group_size = query_shape[i].get_length() / key_shape[i].get_length();
|
||||
}
|
||||
}
|
||||
|
|
@ -62,44 +216,73 @@ struct scaled_dot_product_attention_impl : typed_primitive_impl_ocl<scaled_dot_p
|
|||
if (query_shape[query_shape.size() - 1].is_static())
|
||||
config.head_size = query_shape[query_shape.size() - 1].get_length();
|
||||
|
||||
config.is_causal = prim->is_causal;
|
||||
config.is_causal = desc->is_causal;
|
||||
|
||||
return config;
|
||||
}
|
||||
|
||||
static kernel_params_t get_kernel_params(const kernel_impl_params& impl_param, bool is_dynamic) {
|
||||
public:
|
||||
static kernel_params_t get_kernel_params(const kernel_impl_params& impl_param, bool is_dynamic, bool indirect = false) {
|
||||
const auto& desc = impl_param.typed_desc<scaled_dot_product_attention>();
|
||||
auto params = get_default_params<kernel_selector::sdpa_params>(impl_param, is_dynamic);
|
||||
|
||||
const auto inputs_num = impl_param.input_layouts.size();
|
||||
params.inputs.resize(inputs_num);
|
||||
for (size_t i = 0; i < inputs_num; i++) {
|
||||
auto data_inputs_num = impl_param.input_layouts.size();
|
||||
if (has_indirect_inputs(impl_param))
|
||||
data_inputs_num--;
|
||||
|
||||
params.inputs.resize(data_inputs_num);
|
||||
for (size_t i = 0; i < data_inputs_num; i++) {
|
||||
params.inputs[i] = convert_data_tensor(impl_param.get_input_layout(i));
|
||||
}
|
||||
|
||||
params.conf = get_sdpa_configuration(impl_param);
|
||||
|
||||
const auto& prim = impl_param.typed_desc<scaled_dot_product_attention>();
|
||||
params.input0_order = prim->input_q_transpose_order;
|
||||
params.input1_order = prim->input_k_transpose_order;
|
||||
params.input2_order = prim->input_v_transpose_order;
|
||||
params.output_order = prim->output_transpose_order;
|
||||
params.input0_order = desc->input_q_transpose_order;
|
||||
params.input1_order = desc->input_k_transpose_order;
|
||||
params.input2_order = desc->input_v_transpose_order;
|
||||
params.output_order = desc->output_transpose_order;
|
||||
|
||||
if (indirect && has_indirect_inputs(impl_param)) {
|
||||
params.beam_table = convert_data_tensor(impl_param.get_input_layout(get_beam_table_id(desc)));
|
||||
params.indirect_axis = desc->indirect_axis;
|
||||
}
|
||||
|
||||
params.set_dynamic_shape_offsets();
|
||||
|
||||
// Need to adjust sdpa kernel offset to consider beam table input
|
||||
if (has_indirect_inputs(impl_param)) {
|
||||
auto out_offset = params.outputs[0].get_dynamic_shape_offset();
|
||||
if (indirect)
|
||||
params.beam_table.SetDynamicShapeOffset(out_offset);
|
||||
|
||||
params.outputs[0].SetDynamicShapeOffset(out_offset + kernel_selector::DataTensor::max_rank());
|
||||
}
|
||||
|
||||
return params;
|
||||
}
|
||||
|
||||
static std::unique_ptr<primitive_impl> create(const typed_program_node<scaled_dot_product_attention>& arg, const kernel_impl_params& impl_param) {
|
||||
std::vector<kernel_selector::kernel_data> kernels_data;
|
||||
auto sdpa_kernel_params = get_kernel_params(impl_param, impl_param.is_dynamic());
|
||||
auto& sdpa_kernel_selector = kernel_selector_t::Instance();
|
||||
auto kd = sdpa_kernel_selector.get_best_kernel(sdpa_kernel_params);
|
||||
auto& kernel_selector = kernel_selector_t::Instance();
|
||||
kernels_data.push_back(kernel_selector.get_best_kernel(sdpa_kernel_params));
|
||||
|
||||
return cldnn::make_unique<scaled_dot_product_attention_impl>(kd);
|
||||
if (has_indirect_inputs(impl_param)) {
|
||||
auto indirect_kernel_params = get_kernel_params(impl_param, impl_param.is_dynamic(), true);
|
||||
kernels_data.push_back(kernel_selector.get_best_kernel(indirect_kernel_params));
|
||||
}
|
||||
|
||||
return cldnn::make_unique<scaled_dot_product_attention_impl>(kernels_data);
|
||||
}
|
||||
|
||||
void update_dispatch_data(const kernel_impl_params& impl_param) override {
|
||||
auto kernel_params = get_kernel_params(impl_param, true);
|
||||
(_kernel_data.update_dispatch_data_func)(kernel_params, _kernel_data);
|
||||
(_kernels_data[default_sdpa].update_dispatch_data_func)(kernel_params, _kernels_data[default_sdpa]);
|
||||
|
||||
if (_kernels_data.size() == 2) {
|
||||
auto kernel_params = get_kernel_params(impl_param, true);
|
||||
(_kernels_data[indirect_sdpa].update_dispatch_data_func)(kernel_params, _kernels_data[indirect_sdpa]);
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
|
|
|
|||
|
|
@ -32,6 +32,9 @@ public:
|
|||
static std::vector<layout> calc_output_layouts(scaled_dot_product_attention_node const& /*node*/, const kernel_impl_params& impl_param);
|
||||
static layout calc_output_layout(scaled_dot_product_attention_node const& node, kernel_impl_params const& impl_param);
|
||||
static std::string to_string(scaled_dot_product_attention_node const& node);
|
||||
bool has_indirect_inputs() const {
|
||||
return get_typed_desc<scaled_dot_product_attention>()->indirect_axis != -1;
|
||||
}
|
||||
|
||||
typed_primitive_inst(network& network, scaled_dot_product_attention_node const& desc);
|
||||
};
|
||||
|
|
|
|||
|
|
@ -96,6 +96,24 @@ inline uint FUNC(get_input2_index)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint
|
|||
#endif
|
||||
}
|
||||
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
inline uint FUNC(get_bt_index_nt)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w, uint z, uint y, uint x) {
|
||||
#if BEAM_TABLE_SIMPLE
|
||||
return GET_DATA_INDEX_6D_SAFE(BEAM_TABLE, b, f, w, z, y, x);
|
||||
#else
|
||||
# error sdpa_ref.cl : Unsupported beam table format
|
||||
#endif
|
||||
}
|
||||
|
||||
inline uint FUNC(get_bt_index_key)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w, uint z, uint y, uint x) {
|
||||
return FUNC_CALL(get_bt_index_nt)(OPTIONAL_SHAPE_INFO_TENSOR INPUT1_DIMS_ORDER);
|
||||
}
|
||||
|
||||
inline uint FUNC(get_bt_index_value)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w, uint z, uint y, uint x) {
|
||||
return FUNC_CALL(get_bt_index_nt)(OPTIONAL_SHAPE_INFO_TENSOR INPUT2_DIMS_ORDER);
|
||||
}
|
||||
#endif
|
||||
|
||||
#define VALUE_BLOCK_READ(ptr, offset) BLOCK_READN(INPUT2_TYPE, 1, ptr, offset)
|
||||
#define SUBGROUPS_PER_WG (HEAD_SIZE / SUBGROUP_SIZE)
|
||||
|
||||
|
|
@ -117,6 +135,9 @@ KERNEL(sdpa_opt)(
|
|||
const __global INPUT4_TYPE* scale,
|
||||
#endif
|
||||
__global OUTPUT_TYPE* output,
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
const __global BEAM_TABLE_TYPE* beam_table,
|
||||
#endif
|
||||
__global SOFTMAX_ACCUMULATOR_TYPE* exp_sums,
|
||||
__global SOFTMAX_ACCUMULATOR_TYPE* max_logits,
|
||||
__global OUTPUT_TYPE* tmp_out
|
||||
|
|
@ -125,12 +146,7 @@ KERNEL(sdpa_opt)(
|
|||
const uint batch_idx = get_global_id(0);
|
||||
const uint b0_idx = batch_idx / NUM_HEADS; /* BATCH dim */
|
||||
const uint b1_idx = batch_idx % NUM_HEADS; /* HEADS_NUM dim */
|
||||
|
||||
#if TARGET_SEQ_LEN_BLOCK_SIZE > 1
|
||||
const uint target_seq_idx = (uint)get_global_id(1) * TARGET_SEQ_LEN_BLOCK_SIZE;
|
||||
#else
|
||||
const uint target_seq_idx = get_global_id(1);
|
||||
#endif
|
||||
const uint lid = get_local_id(2);
|
||||
const uint head_size_idx = lid;
|
||||
|
||||
|
|
@ -173,12 +189,7 @@ KERNEL(sdpa_opt)(
|
|||
// Query input loading to SLM
|
||||
#define QUERY_STEP_LOCAL SUBGROUP_SIZE * SUBGROUPS_PER_WG
|
||||
uint query_local_offset = sgid * SUBGROUP_SIZE + sglid;
|
||||
|
||||
#if TARGET_SEQ_LEN_BLOCK_SIZE > 1
|
||||
const uint seq_idx_end = min(TARGET_SEQ_LEN - target_seq_idx, (uint)TARGET_SEQ_LEN_BLOCK_SIZE);
|
||||
#else
|
||||
const uint seq_idx_end = 1;
|
||||
#endif
|
||||
#ifdef INPUT0_DIMS_ORDER
|
||||
uint query_offset = FUNC_CALL(get_input0_index)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, target_seq_idx, (sgid * SUBGROUP_SIZE));
|
||||
uint query_offset_next_seq = FUNC_CALL(get_input0_index)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, target_seq_idx + 1, (sgid * SUBGROUP_SIZE));
|
||||
|
|
@ -207,9 +218,14 @@ KERNEL(sdpa_opt)(
|
|||
// HEAD_SIZE / SUBGROUPS_PER_WG times in the loop and saves the result to the qk_local SLM buffer
|
||||
for (uint seq_len = sgid; seq_len < partition_seq_len; seq_len += (HEAD_SIZE / SUBGROUP_SIZE)) {
|
||||
#ifdef INPUT1_DIMS_ORDER
|
||||
uint key_offset = FUNC_CALL(get_input1_index)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + seq_len, 0);
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
const uint b_idx = beam_table[FUNC_CALL(get_bt_index_key)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + seq_len, 0)];
|
||||
#else
|
||||
uint key_offset = INPUT1_GET_INDEX(b0_idx, b1_idx, start_partition_idx + seq_len, 0);
|
||||
const uint b_idx = b0_idx;
|
||||
#endif
|
||||
const uint key_offset = FUNC_CALL(get_input1_index)(OPTIONAL_SHAPE_INFO_TENSOR b_idx, b1_idx, 0, 0, start_partition_idx + seq_len, 0);
|
||||
#else
|
||||
const uint key_offset = INPUT1_GET_INDEX(b0_idx, b1_idx, start_partition_idx + seq_len, 0);
|
||||
#endif
|
||||
|
||||
INPUT0_TYPE acc[TARGET_SEQ_LEN_BLOCK_SIZE] = {INPUT0_VAL_ZERO};
|
||||
|
|
@ -316,11 +332,7 @@ KERNEL(sdpa_opt)(
|
|||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
INPUT0_TYPE qk_val[TARGET_SEQ_LEN_BLOCK_SIZE];
|
||||
#if TARGET_SEQ_LEN_BLOCK_SIZE > 1
|
||||
const uint seq_idx_end = min(TARGET_SEQ_LEN - target_seq_idx, (uint)TARGET_SEQ_LEN_BLOCK_SIZE);
|
||||
#else
|
||||
const uint seq_idx_end = 1;
|
||||
#endif
|
||||
for (uint seq_idx = 0; seq_idx < seq_idx_end; seq_idx++) {
|
||||
// Iterate over all values QK values in SLM and apply scale and attention mask
|
||||
for (uint seq_len = sgid * SUBGROUP_SIZE + sglid; seq_len < partition_seq_len; seq_len += (HEAD_SIZE)) {
|
||||
|
|
@ -349,11 +361,7 @@ KERNEL(sdpa_opt)(
|
|||
|
||||
{
|
||||
// SoftMax calculation
|
||||
#if TARGET_SEQ_LEN_BLOCK_SIZE > 1
|
||||
const uint seq_idx_end = min(TARGET_SEQ_LEN - target_seq_idx, (uint)TARGET_SEQ_LEN_BLOCK_SIZE);
|
||||
#else
|
||||
const uint seq_idx_end = 1;
|
||||
#endif
|
||||
// Find the maximum value of qk in the subgroup
|
||||
for (uint seq_idx = 0; seq_idx < seq_idx_end; seq_idx++) {
|
||||
qk_max[seq_idx] = sub_group_reduce_max(qk_max[seq_idx]);
|
||||
|
|
@ -446,20 +454,26 @@ KERNEL(sdpa_opt)(
|
|||
{
|
||||
// Gemm2 calculation
|
||||
OUTPUT_TYPE acc[TARGET_SEQ_LEN_BLOCK_SIZE] = {OUTPUT_VAL_ZERO};
|
||||
|
||||
#ifndef BEAM_TABLE_TYPE
|
||||
#ifdef INPUT2_DIMS_ORDER
|
||||
uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, 0, 0);
|
||||
uint value_offset_next_seq = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, 1, 0);
|
||||
const uint value_pitch = value_offset_next_seq - value_offset;
|
||||
#else
|
||||
const uint value_pitch = HEAD_SIZE;
|
||||
#endif
|
||||
#endif
|
||||
|
||||
for (uint seq_len = 0; seq_len < partition_seq_len / SUBGROUP_SIZE; seq_len++) {
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
uint b_idx = beam_table[FUNC_CALL(get_bt_index_value)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + (seq_len * SUBGROUP_SIZE) + sglid, (head_size_idx / SUBGROUP_SIZE) * SUBGROUP_SIZE)];
|
||||
uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b_idx, b1_idx, 0, 0, start_partition_idx + (seq_len * SUBGROUP_SIZE) + sglid, (head_size_idx / SUBGROUP_SIZE) * SUBGROUP_SIZE);
|
||||
#else
|
||||
#ifdef INPUT2_DIMS_ORDER
|
||||
uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + (seq_len * SUBGROUP_SIZE), head_size_idx);
|
||||
#else
|
||||
uint value_offset = INPUT2_GET_INDEX(b0_idx, b1_idx, start_partition_idx + (seq_len * SUBGROUP_SIZE), head_size_idx);
|
||||
#endif
|
||||
#endif
|
||||
|
||||
OUTPUT_TYPE qk_val[TARGET_SEQ_LEN_BLOCK_SIZE];
|
||||
|
|
@ -468,19 +482,30 @@ KERNEL(sdpa_opt)(
|
|||
}
|
||||
|
||||
unroll_for (uint i = 0; i < SUBGROUP_SIZE; i++) {
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
INPUT2_TYPE value_val = VALUE_BLOCK_READ(value_input, sub_group_broadcast(value_offset, i));
|
||||
#else
|
||||
INPUT2_TYPE value_val = VALUE_BLOCK_READ(value_input, value_offset);
|
||||
#endif
|
||||
unroll_for (uint seq_idx = 0; seq_idx < TARGET_SEQ_LEN_BLOCK_SIZE; seq_idx++) {
|
||||
acc[seq_idx] = mad(sub_group_broadcast(qk_val[seq_idx], i), value_val, acc[seq_idx]);
|
||||
}
|
||||
|
||||
#ifndef BEAM_TABLE_TYPE
|
||||
value_offset += value_pitch;
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
const uint seq_len_leftovers_start = (partition_seq_len / SUBGROUP_SIZE) * SUBGROUP_SIZE;
|
||||
for (uint seq_len = seq_len_leftovers_start; seq_len < partition_seq_len; seq_len++) {
|
||||
#ifdef INPUT2_DIMS_ORDER
|
||||
const uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + seq_len, head_size_idx);
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
const uint b_idx = beam_table[FUNC_CALL(get_bt_index_value)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + seq_len, head_size_idx)];
|
||||
#else
|
||||
const uint b_idx = b0_idx;
|
||||
#endif
|
||||
const uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b_idx, b1_idx, 0, 0, start_partition_idx + seq_len, head_size_idx);
|
||||
#else
|
||||
const uint value_offset = INPUT2_GET_INDEX(b0_idx, b1_idx, start_partition_idx + seq_len, head_size_idx);
|
||||
#endif
|
||||
|
|
@ -500,11 +525,7 @@ KERNEL(sdpa_opt)(
|
|||
// If the number of partitions is greater than 1, save results to the temporary buffer;
|
||||
// otherwise, save results directly to the main output.
|
||||
if (num_of_partitions > 1) {
|
||||
#if TARGET_SEQ_LEN_BLOCK_SIZE > 1
|
||||
const uint seq_idx_end = min(TARGET_SEQ_LEN - target_seq_idx, (uint)TARGET_SEQ_LEN_BLOCK_SIZE);
|
||||
#else
|
||||
const uint seq_idx_end = 1;
|
||||
#endif
|
||||
for (uint seq_idx = 0; seq_idx < seq_idx_end; seq_idx++) {
|
||||
// Data layout of tmp_output buf: [batch, heads_num, q_len, partition_idx, head_size]
|
||||
const uint tmp_out_offset = b0_idx * (NUM_HEADS * TARGET_SEQ_LEN * num_of_partitions * HEAD_SIZE) +
|
||||
|
|
@ -515,15 +536,11 @@ KERNEL(sdpa_opt)(
|
|||
tmp_out[tmp_out_offset] = acc[seq_idx];
|
||||
}
|
||||
} else {
|
||||
#if TARGET_SEQ_LEN_BLOCK_SIZE > 1
|
||||
const uint seq_idx_end = min(TARGET_SEQ_LEN - target_seq_idx, (uint)TARGET_SEQ_LEN_BLOCK_SIZE);
|
||||
#else
|
||||
const uint seq_idx_end = 1;
|
||||
#endif
|
||||
for (uint seq_idx = 0; seq_idx < seq_idx_end; seq_idx++) {
|
||||
const uint output_offset = OUTPUT_GET_INDEX(b0_idx, b1_idx, target_seq_idx + seq_idx, head_size_idx);
|
||||
const uint output_offset = OUTPUT_GET_INDEX(b0_idx, b1_idx, target_seq_idx + seq_idx, head_size_idx);
|
||||
|
||||
output[output_offset] = acc[seq_idx];
|
||||
output[output_offset] = acc[seq_idx];
|
||||
}
|
||||
}
|
||||
} // Gemm2 calculation end
|
||||
|
|
@ -545,6 +562,9 @@ KERNEL(sdpa_opt)(
|
|||
const __global INPUT4_TYPE* scale,
|
||||
#endif
|
||||
__global OUTPUT_TYPE* output,
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
const __global BEAM_TABLE_TYPE* beam_table,
|
||||
#endif
|
||||
__global SOFTMAX_ACCUMULATOR_TYPE* exp_sums,
|
||||
__global SOFTMAX_ACCUMULATOR_TYPE* max_logits,
|
||||
__global OUTPUT_TYPE* tmp_out
|
||||
|
|
@ -637,6 +657,10 @@ KERNEL(sdpa_opt)(
|
|||
// Main Gemm1 calculation loop
|
||||
uint seq_len = sgid * TARGET_SEQ_LEN_BLOCK_SIZE;
|
||||
for (; seq_len < partition_seq_len; seq_len += SUBGROUPS_PER_WG * SUBGROUP_SIZE) {
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
const uint b_idx = beam_table[FUNC_CALL(get_bt_index_key)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + seq_len + sglid, 0)];
|
||||
const uint key_offset = FUNC_CALL(get_input1_index)(OPTIONAL_SHAPE_INFO_TENSOR b_idx, b1_idx, 0, 0, start_partition_idx + seq_len + sglid, 0);
|
||||
#else
|
||||
#ifdef INPUT1_DIMS_ORDER
|
||||
uint key_offset = FUNC_CALL(get_input1_index)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + seq_len, 0);
|
||||
uint key_offset_next_seq = FUNC_CALL(get_input1_index)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + seq_len + 1, 0);
|
||||
|
|
@ -644,6 +668,7 @@ KERNEL(sdpa_opt)(
|
|||
#else
|
||||
uint key_offset = INPUT1_GET_INDEX(b0_idx, b1_idx, start_partition_idx + seq_len, 0);
|
||||
const uint key_pitch = HEAD_SIZE;
|
||||
#endif
|
||||
#endif
|
||||
|
||||
INPUT0_TYPE acc[TARGET_SEQ_LEN_BLOCK_SIZE] = {INPUT0_VAL_ZERO};
|
||||
|
|
@ -660,7 +685,11 @@ KERNEL(sdpa_opt)(
|
|||
}
|
||||
|
||||
unroll_for (uint key_row_idx = 0; key_row_idx < TARGET_SEQ_LEN_BLOCK_SIZE; key_row_idx++) {
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
INPUT1_TYPE key_vals = KEY_BLOCK_READ(key_input, sub_group_broadcast(key_offset, key_row_idx) + head_idx_index);
|
||||
#else
|
||||
INPUT1_TYPE key_vals = KEY_BLOCK_READ(key_input, key_offset + key_row_idx * key_pitch + head_idx_index);
|
||||
#endif
|
||||
|
||||
unroll_for (uint i = 0; i < SUBGROUP_SIZE; i++) {
|
||||
acc[key_row_idx] = mad(sub_group_broadcast(key_vals, i), queries_vec[i], acc[key_row_idx]);
|
||||
|
|
@ -801,10 +830,15 @@ KERNEL(sdpa_opt)(
|
|||
#endif
|
||||
|
||||
for (uint seq_len = 0; seq_len < partition_seq_len / SUBGROUP_SIZE; seq_len++) {
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
const uint b_idx = beam_table[FUNC_CALL(get_bt_index_value)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + (seq_len * SUBGROUP_SIZE) + sglid, sgid * SUBGROUP_SIZE)];
|
||||
const uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b_idx, b1_idx, 0, 0, start_partition_idx + (seq_len * SUBGROUP_SIZE) + sglid, sgid * SUBGROUP_SIZE);
|
||||
#else
|
||||
#ifdef INPUT2_DIMS_ORDER
|
||||
uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + (seq_len * SUBGROUP_SIZE), head_size_idx);
|
||||
#else
|
||||
uint value_offset = INPUT2_GET_INDEX(b0_idx, b1_idx, start_partition_idx + (seq_len * SUBGROUP_SIZE), head_size_idx);
|
||||
#endif
|
||||
#endif
|
||||
|
||||
OUTPUT_TYPE qk_val[TARGET_SEQ_LEN_BLOCK_SIZE];
|
||||
|
|
@ -813,12 +847,18 @@ KERNEL(sdpa_opt)(
|
|||
}
|
||||
|
||||
unroll_for (uint i = 0; i < SUBGROUP_SIZE; i++) {
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
INPUT2_TYPE value_val = VALUE_BLOCK_READ(value_input, sub_group_broadcast(value_offset, i));
|
||||
#else
|
||||
INPUT2_TYPE value_val = VALUE_BLOCK_READ(value_input, value_offset);
|
||||
#endif
|
||||
unroll_for (uint seq_idx = 0; seq_idx < TARGET_SEQ_LEN_BLOCK_SIZE; seq_idx++) {
|
||||
acc[seq_idx] = mad(sub_group_broadcast(qk_val[seq_idx], i), value_val, acc[seq_idx]);
|
||||
}
|
||||
|
||||
#ifndef BEAM_TABLE_TYPE
|
||||
value_offset += value_pitch;
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -833,17 +873,31 @@ KERNEL(sdpa_opt)(
|
|||
qk_val[seq_idx] = qk_local[qk_offset];
|
||||
qk_offset += SEQ_LEN_PARTITION_SIZE;
|
||||
}
|
||||
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
const uint b_idx = beam_table[FUNC_CALL(get_bt_index_value)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + seq_len_leftovers_start + sglid, sgid * SUBGROUP_SIZE)];
|
||||
const uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b_idx, b1_idx, 0, 0, start_partition_idx + seq_len_leftovers_start + sglid, sgid * SUBGROUP_SIZE);
|
||||
#else
|
||||
#ifdef INPUT2_DIMS_ORDER
|
||||
uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + seq_len_leftovers_start, head_size_idx);
|
||||
#else
|
||||
uint value_offset = INPUT2_GET_INDEX(b0_idx, b1_idx, start_partition_idx + seq_len_leftovers_start, head_size_idx);
|
||||
#endif
|
||||
#endif
|
||||
|
||||
for (uint seq_len_idx = 0; seq_len_idx < partition_seq_len - seq_len_leftovers_start; seq_len_idx++) {
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
INPUT2_TYPE value_val = VALUE_BLOCK_READ(value_input, sub_group_broadcast(value_offset, seq_len_idx));
|
||||
#else
|
||||
INPUT2_TYPE value_val = VALUE_BLOCK_READ(value_input, value_offset);
|
||||
#endif
|
||||
|
||||
for (uint seq_idx = 0; seq_idx < TARGET_SEQ_LEN_BLOCK_SIZE; seq_idx++) {
|
||||
acc[seq_idx] = mad(sub_group_broadcast(qk_val[seq_idx], seq_len_idx), value_val, acc[seq_idx]);
|
||||
}
|
||||
|
||||
#ifndef BEAM_TABLE_TYPE
|
||||
value_offset += value_pitch;
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -890,10 +944,15 @@ KERNEL(sdpa_opt)(
|
|||
#endif
|
||||
|
||||
for (uint seq_len = 0; seq_len < partition_seq_len / SUBGROUP_SIZE; seq_len++) {
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
const uint b_idx = beam_table[FUNC_CALL(get_bt_index_value)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + (seq_len * SUBGROUP_SIZE) + sglid, sgid * SUBGROUP_SIZE)];
|
||||
const uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b_idx, b1_idx, 0, 0, start_partition_idx + (seq_len * SUBGROUP_SIZE) + sglid, sgid * SUBGROUP_SIZE);
|
||||
#else
|
||||
#ifdef INPUT2_DIMS_ORDER
|
||||
uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b0_idx, b1_idx, 0, 0, start_partition_idx + (seq_len * SUBGROUP_SIZE), head_size_idx);
|
||||
#else
|
||||
uint value_offset = INPUT2_GET_INDEX(b0_idx, b1_idx, start_partition_idx + (seq_len * SUBGROUP_SIZE), head_size_idx);
|
||||
#endif
|
||||
#endif
|
||||
|
||||
OUTPUT_TYPE qk_val[TARGET_SEQ_LEN_BLOCK_SIZE];
|
||||
|
|
@ -902,12 +961,18 @@ KERNEL(sdpa_opt)(
|
|||
}
|
||||
|
||||
unroll_for (uint i = 0; i < SUBGROUP_SIZE; i++) {
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
INPUT2_TYPE value_val = VALUE_BLOCK_READ(value_input, sub_group_broadcast(value_offset, i));
|
||||
#else
|
||||
INPUT2_TYPE value_val = VALUE_BLOCK_READ(value_input, value_offset);
|
||||
#endif
|
||||
unroll_for (uint seq_idx = 0; seq_idx < TARGET_SEQ_LEN_BLOCK_SIZE; seq_idx++) {
|
||||
acc[seq_idx] = mad(sub_group_broadcast(qk_val[seq_idx], i), value_val, acc[seq_idx]);
|
||||
}
|
||||
|
||||
#ifndef BEAM_TABLE_TYPE
|
||||
value_offset += value_pitch;
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -93,6 +93,24 @@ inline uint FUNC(get_input2_index)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint
|
|||
#endif
|
||||
}
|
||||
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
inline uint FUNC(get_bt_index_nt)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w, uint z, uint y, uint x) {
|
||||
#if BEAM_TABLE_SIMPLE
|
||||
return GET_DATA_INDEX_6D_SAFE(BEAM_TABLE, b, f, w, z, y, x);
|
||||
#else
|
||||
# error sdpa_ref.cl : Unsupported beam table format
|
||||
#endif
|
||||
}
|
||||
|
||||
inline uint FUNC(get_bt_index_key)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w, uint z, uint y, uint x) {
|
||||
return FUNC_CALL(get_bt_index_nt)(OPTIONAL_SHAPE_INFO_TENSOR INPUT1_DIMS_ORDER);
|
||||
}
|
||||
|
||||
inline uint FUNC(get_bt_index_value)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w, uint z, uint y, uint x) {
|
||||
return FUNC_CALL(get_bt_index_nt)(OPTIONAL_SHAPE_INFO_TENSOR INPUT2_DIMS_ORDER);
|
||||
}
|
||||
#endif
|
||||
|
||||
#define APPLY_SCALE_TO_QUERY 1
|
||||
|
||||
KERNEL(sdpa_ref)(
|
||||
|
|
@ -107,6 +125,9 @@ KERNEL(sdpa_ref)(
|
|||
const __global INPUT4_TYPE* scale,
|
||||
#endif
|
||||
__global OUTPUT_TYPE* output,
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
const __global BEAM_TABLE_TYPE* beam_table,
|
||||
#endif
|
||||
__global OUTPUT_TYPE* tmp_buf
|
||||
)
|
||||
{
|
||||
|
|
@ -129,7 +150,12 @@ KERNEL(sdpa_ref)(
|
|||
OUTPUT_TYPE acc = 0;
|
||||
for (uint h = 0; h < HEAD_SIZE /* head_size */; h++) {
|
||||
uint query_offset = FUNC_CALL(get_input0_index)(OPTIONAL_SHAPE_INFO_TENSOR b0, b1, 0, 0, target_seq_idx, h);
|
||||
uint key_offset = FUNC_CALL(get_input1_index)(OPTIONAL_SHAPE_INFO_TENSOR b0, b1, 0, 0, s, h);
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
uint b_idx = beam_table[FUNC_CALL(get_bt_index_key)(OPTIONAL_SHAPE_INFO_TENSOR b0, b1, 0, 0, s, h)];
|
||||
#else
|
||||
uint b_idx = b0;
|
||||
#endif
|
||||
uint key_offset = FUNC_CALL(get_input1_index)(OPTIONAL_SHAPE_INFO_TENSOR b_idx, b1, 0, 0, s, h);
|
||||
|
||||
#if APPLY_SCALE_TO_QUERY
|
||||
INPUT0_TYPE q_val = query_input[query_offset] * scale_val;
|
||||
|
|
@ -202,7 +228,13 @@ KERNEL(sdpa_ref)(
|
|||
uint tmp_buf_offset = b0 * (NUM_HEADS * TARGET_SEQ_LEN * SOURCE_SEQ_LEN) +
|
||||
b1 * (TARGET_SEQ_LEN * SOURCE_SEQ_LEN) +
|
||||
target_seq_idx * (SOURCE_SEQ_LEN) + s;
|
||||
uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b0, b1, 0, 0, s, head_size_idx);
|
||||
|
||||
#ifdef BEAM_TABLE_TYPE
|
||||
uint b_idx = beam_table[FUNC_CALL(get_bt_index_value)(OPTIONAL_SHAPE_INFO_TENSOR b0, b1, 0, 0, s, head_size_idx)];
|
||||
#else
|
||||
uint b_idx = b0;
|
||||
#endif
|
||||
uint value_offset = FUNC_CALL(get_input2_index)(OPTIONAL_SHAPE_INFO_TENSOR b_idx, b1, 0, 0, s, head_size_idx);
|
||||
|
||||
acc += tmp_buf[tmp_buf_offset] * value_input[value_offset];
|
||||
}
|
||||
|
|
|
|||
|
|
@ -155,7 +155,7 @@ FullyConnectedKernelIMAD::FullyConnectedTuningData FullyConnectedKernelIMAD::Get
|
|||
}
|
||||
|
||||
// In most cases SIMD8 works faster than SIMD16
|
||||
tuning_data.sub_group_size = 8;
|
||||
tuning_data.sub_group_size = IsSIMDSizeSupported(params.engineInfo, 8) ? 8 : 16;
|
||||
|
||||
if (!params.is_shape_agnostic) {
|
||||
auto mk_size = if_num * ib_num;
|
||||
|
|
|
|||
|
|
@ -85,15 +85,30 @@ JitConstants SDPAKernelBase::GetJitConstants(const sdpa_params& params) const {
|
|||
return true;
|
||||
};
|
||||
|
||||
if ((!params.input0_order.empty() && !is_default_order(params.input0_order)) || params.conf.broadcast_axis != -1) {
|
||||
auto use_index_calc_func = [&](const std::vector<int64_t> order, bool is_query = false) {
|
||||
if (!params.input0_order.empty() && !is_default_order(params.input0_order))
|
||||
return true;
|
||||
|
||||
if (params.conf.broadcast_axis != -1)
|
||||
return true;
|
||||
|
||||
if (params.indirect_axis != -1 && !is_query)
|
||||
return true;
|
||||
|
||||
return false;
|
||||
};
|
||||
|
||||
if (params.indirect_axis != -1)
|
||||
jit.AddConstant(MakeJitConstant("BEAM_TABLE", params.beam_table));
|
||||
|
||||
if (use_index_calc_func(params.input0_order, true))
|
||||
jit.AddConstant(MakeJitConstant("INPUT0_DIMS_ORDER", GetDimsOrder(params.input0_order)));
|
||||
}
|
||||
if ((!params.input1_order.empty() && !is_default_order(params.input1_order)) || params.conf.broadcast_axis != -1) {
|
||||
|
||||
if (use_index_calc_func(params.input1_order))
|
||||
jit.AddConstant(MakeJitConstant("INPUT1_DIMS_ORDER", GetDimsOrder(params.input1_order)));
|
||||
}
|
||||
if ((!params.input2_order.empty() && !is_default_order(params.input2_order)) || params.conf.broadcast_axis != -1) {
|
||||
|
||||
if (use_index_calc_func(params.input2_order))
|
||||
jit.AddConstant(MakeJitConstant("INPUT2_DIMS_ORDER", GetDimsOrder(params.input2_order)));
|
||||
}
|
||||
|
||||
TransposedDimensionAccessHelperJit dims_q(params.inputs[0], params.input0_order);
|
||||
jit.AddConstant(MakeJitConstant("TARGET_SEQ_LEN", dims_q.y()));
|
||||
|
|
|
|||
|
|
@ -99,6 +99,9 @@ struct sdpa_params : public base_params {
|
|||
std::vector<int64_t> input1_order;
|
||||
std::vector<int64_t> input2_order;
|
||||
std::vector<int64_t> output_order;
|
||||
int64_t indirect_axis = -1;
|
||||
|
||||
DataTensor beam_table;
|
||||
|
||||
sdpa_configuration conf;
|
||||
};
|
||||
|
|
|
|||
|
|
@ -29,6 +29,22 @@ static size_t get_seq_len_partition_size() {
|
|||
return seq_len;
|
||||
}
|
||||
|
||||
static std::string GetKernelName(std::string base_name, KernelsTypes type, bool is_indirect) {
|
||||
auto kernel_name = base_name;
|
||||
if (is_indirect)
|
||||
kernel_name += "_ind";
|
||||
|
||||
if (type == KernelsTypes::SINGLE_TOKEN) {
|
||||
kernel_name += "_single_token";
|
||||
} else if (type == KernelsTypes::MULTI_TOKENS) {
|
||||
kernel_name += "_multi_tokens";
|
||||
} else if (type == KernelsTypes::FINALIZATION) {
|
||||
kernel_name += "_finalization";
|
||||
}
|
||||
|
||||
return kernel_name;
|
||||
}
|
||||
|
||||
ParamsKey SDPAKernelOpt::GetSupportedKey() const {
|
||||
ParamsKey k;
|
||||
k.EnableInputDataType(Datatype::F16);
|
||||
|
|
@ -104,7 +120,7 @@ CommonDispatchData SDPAKernelOpt::SetDefault(const sdpa_params& params, size_t k
|
|||
CeilDiv(target_seq_len, target_seq_len_block_size),
|
||||
head_size * num_of_partitions };
|
||||
dispatch_data.lws = { 1, 1, head_size };
|
||||
} else if (kernel_idx == 2) {
|
||||
} else if (kernel_idx == KernelsTypes::FINALIZATION) {
|
||||
dispatch_data.gws = { batch_size * heads_num,
|
||||
target_seq_len,
|
||||
16 };
|
||||
|
|
@ -134,8 +150,7 @@ KernelsData SDPAKernelOpt::GetKernelsData(const Params& params) const {
|
|||
const auto& prim_params = dynamic_cast<const sdpa_params&>(params);
|
||||
for (size_t kernel_idx = 0; kernel_idx < kernels_num; kernel_idx++) {
|
||||
auto dispatch_data = SetDefault(prim_params, kernel_idx);
|
||||
auto kernel_name = kernel_idx == 0 ? kernelName + "_single_token" :
|
||||
kernel_idx == 1 ? kernelName + "_multi_tokens" : kernelName + "_finalization";
|
||||
auto kernel_name = GetKernelName(kernelName, static_cast<KernelsTypes>(kernel_idx), prim_params.indirect_axis != -1);
|
||||
auto entry_point = GetEntryPoint(kernel_name, prim_params.layerID, params);
|
||||
auto jit_constants = GetJitConstants(prim_params, kernel_idx);
|
||||
auto jit = CreateJit(kernel_name, jit_constants, entry_point);
|
||||
|
|
@ -171,6 +186,9 @@ KernelsData SDPAKernelOpt::GetKernelsData(const Params& params) const {
|
|||
auto tmp_out_elements_count = (num_of_partitions == 1) ? 1 : output.LogicalSize() * num_of_partitions;
|
||||
auto tmp_out_size = tmp_out_elements_count * tmp_out_dt_size;
|
||||
|
||||
if (prim_params.indirect_axis != -1 && kernel_idx != KernelsTypes::FINALIZATION)
|
||||
kernel.params.arguments.push_back({ArgumentDescriptor::Types::INPUT, static_cast<uint32_t>(prim_params.inputs.size())});
|
||||
|
||||
kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 0});
|
||||
kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 1});
|
||||
kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 2});
|
||||
|
|
|
|||
|
|
@ -13,6 +13,8 @@ ParamsKey SDPAKernelRef::GetSupportedKey() const {
|
|||
ParamsKey k;
|
||||
k.EnableInputDataType(Datatype::F16);
|
||||
k.EnableInputDataType(Datatype::F32);
|
||||
// beam table input
|
||||
k.EnableInputDataType(Datatype::INT32);
|
||||
|
||||
k.EnableOutputDataType(Datatype::F16);
|
||||
k.EnableOutputDataType(Datatype::F32);
|
||||
|
|
@ -72,6 +74,9 @@ KernelsData SDPAKernelRef::GetKernelsData(const Params& params) const {
|
|||
"", false, false, static_cast<int>(prim_params.inputs.size()),
|
||||
GetFusedPrimitiveInputsCount(params), 1, prim_params.is_shape_agnostic);
|
||||
|
||||
if (prim_params.indirect_axis != -1)
|
||||
kernel.params.arguments.push_back({ArgumentDescriptor::Types::INPUT, static_cast<uint32_t>(prim_params.inputs.size())});
|
||||
|
||||
kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 0});
|
||||
|
||||
kd.internalBufferSizes.clear();
|
||||
|
|
|
|||
|
|
@ -17,6 +17,7 @@
|
|||
#include "openvino/op/split.hpp"
|
||||
#include "openvino/op/prelu.hpp"
|
||||
#include "openvino/op/roi_align.hpp"
|
||||
#include "openvino/op/roi_align_rotated.hpp"
|
||||
#include "openvino/op/variadic_split.hpp"
|
||||
#include "openvino/op/util/op_types.hpp"
|
||||
#include "openvino/op/loop.hpp"
|
||||
|
|
@ -221,7 +222,8 @@ static void CreateConstantOp(ProgramBuilder& p, const std::shared_ptr<ov::op::v0
|
|||
if (constDims.size() == 4 && input_shape.size() == 3) { // In case of weight dim 4 and input dim 3,
|
||||
constDims.push_back(1); // The weight cldnn tensor adds 1d to the end as the input cldnn tensor does
|
||||
}
|
||||
} else if (ov::is_type<ov::op::v3::ROIAlign>(outOp) || ov::is_type<ov::op::v9::ROIAlign>(outOp)) {
|
||||
} else if (ov::is_type<ov::op::v3::ROIAlign>(outOp) || ov::is_type<ov::op::v9::ROIAlign>(outOp) ||
|
||||
ov::is_type<ov::op::v15::ROIAlignRotated>(outOp)) { //< Hacks...
|
||||
consts[op].needsBatchInterpretation = constDims.size() == 1;
|
||||
} else if ((ov::is_type<ov::op::v5::Loop>(outOp) || ov::is_type<ov::op::v0::TensorIterator>(outOp))) {
|
||||
// when inner network has 1d parameter which is connected to outer loop's constant 1d data,
|
||||
|
|
|
|||
|
|
@ -6,6 +6,7 @@
|
|||
#include "intel_gpu/plugin/common_utils.hpp"
|
||||
|
||||
#include "intel_gpu/op/sdpa.hpp"
|
||||
#include "intel_gpu/op/indirect_sdpa.hpp"
|
||||
|
||||
#include "openvino/op/scaled_dot_product_attention.hpp"
|
||||
|
||||
|
|
@ -15,6 +16,7 @@ namespace ov {
|
|||
namespace op {
|
||||
namespace internal {
|
||||
using SDPA = ov::intel_gpu::op::SDPA;
|
||||
using IndirectSDPA = ov::intel_gpu::op::IndirectSDPA;
|
||||
} // namespace internal
|
||||
} // namespace op
|
||||
} // namespace ov
|
||||
|
|
@ -41,9 +43,30 @@ static void CreateSDPAOp(ProgramBuilder& p, const std::shared_ptr<ov::op::intern
|
|||
auto layerName = layer_type_name_ID(op);
|
||||
|
||||
bool is_causal = op->get_causal();
|
||||
int64_t indirect_axis = -1;
|
||||
auto sdpa_prim = cldnn::scaled_dot_product_attention(layerName,
|
||||
inputs,
|
||||
is_causal,
|
||||
indirect_axis,
|
||||
op->get_input0_transpose_order(),
|
||||
op->get_input1_transpose_order(),
|
||||
op->get_input2_transpose_order(),
|
||||
op->get_output_transpose_order());
|
||||
|
||||
p.add_primitive(*op, sdpa_prim);
|
||||
}
|
||||
|
||||
static void CreateIndirectSDPAOp(ProgramBuilder& p, const std::shared_ptr<ov::op::internal::IndirectSDPA>& op) {
|
||||
validate_inputs_count(op, {4, 5, 6});
|
||||
auto inputs = p.GetInputInfo(op);
|
||||
auto layerName = layer_type_name_ID(op);
|
||||
|
||||
bool is_causal = op->get_causal();
|
||||
int64_t indirect_axis = op->get_indirect_axis();
|
||||
auto sdpa_prim = cldnn::scaled_dot_product_attention(layerName,
|
||||
inputs,
|
||||
is_causal,
|
||||
indirect_axis,
|
||||
op->get_input0_transpose_order(),
|
||||
op->get_input1_transpose_order(),
|
||||
op->get_input2_transpose_order(),
|
||||
|
|
@ -53,6 +76,7 @@ static void CreateSDPAOp(ProgramBuilder& p, const std::shared_ptr<ov::op::intern
|
|||
}
|
||||
|
||||
REGISTER_FACTORY_IMPL(internal, SDPA);
|
||||
REGISTER_FACTORY_IMPL(internal, IndirectSDPA);
|
||||
REGISTER_FACTORY_IMPL(v13, ScaledDotProductAttention);
|
||||
|
||||
} // namespace intel_gpu
|
||||
|
|
|
|||
|
|
@ -542,6 +542,7 @@ std::vector<ov::PropertyName> Plugin::get_supported_properties() const {
|
|||
ov::PropertyName{ov::intel_gpu::hint::host_task_priority.name(), PropertyMutability::RW},
|
||||
ov::PropertyName{ov::intel_gpu::hint::queue_priority.name(), PropertyMutability::RW},
|
||||
ov::PropertyName{ov::intel_gpu::hint::queue_throttle.name(), PropertyMutability::RW},
|
||||
ov::PropertyName{ov::intel_gpu::hint::enable_sdpa_optimization.name(), PropertyMutability::RW},
|
||||
ov::PropertyName{ov::intel_gpu::enable_loop_unrolling.name(), PropertyMutability::RW},
|
||||
ov::PropertyName{ov::intel_gpu::disable_winograd_convolution.name(), PropertyMutability::RW},
|
||||
ov::PropertyName{ov::cache_dir.name(), PropertyMutability::RW},
|
||||
|
|
|
|||
|
|
@ -6,7 +6,9 @@
|
|||
#include <memory>
|
||||
|
||||
#include "intel_gpu/op/gemm.hpp"
|
||||
#include "intel_gpu/op/sdpa.hpp"
|
||||
#include "intel_gpu/op/indirect_gemm.hpp"
|
||||
#include "intel_gpu/op/indirect_sdpa.hpp"
|
||||
#include "intel_gpu/op/kv_cache.hpp"
|
||||
#include "intel_gpu/op/read_value.hpp"
|
||||
#include "intel_gpu/plugin/common_utils.hpp"
|
||||
|
|
@ -42,7 +44,7 @@ void replace_node_unsafe(const std::shared_ptr<ov::Node>& target, const std::sha
|
|||
namespace ov {
|
||||
namespace intel_gpu {
|
||||
|
||||
IndirectKVCache::IndirectKVCache() {
|
||||
IndirectGemmOpt::IndirectGemmOpt() {
|
||||
using namespace ov::pass::pattern;
|
||||
|
||||
auto beam_idx = wrap_type<ov::op::v0::Parameter>();
|
||||
|
|
@ -108,9 +110,141 @@ IndirectKVCache::IndirectKVCache() {
|
|||
return true;
|
||||
};
|
||||
|
||||
auto m = std::make_shared<ov::pass::pattern::Matcher>(matmul, "IndirectKVCache");
|
||||
auto m = std::make_shared<ov::pass::pattern::Matcher>(matmul, "IndirectGemmOpt");
|
||||
this->register_matcher(m, callback);
|
||||
}
|
||||
|
||||
IndirectSDPAOpt::IndirectSDPAOpt() {
|
||||
using namespace ov::pass::pattern;
|
||||
using ov::pass::pattern::op::Or;
|
||||
|
||||
auto beam_idx = wrap_type<ov::op::v0::Parameter>();
|
||||
auto gather_input_0 = wrap_type<ov::intel_gpu::op::ReadValue>();
|
||||
auto gather_input_1 = wrap_type<ov::intel_gpu::op::ReadValue>();
|
||||
auto axis_const = wrap_type<ov::op::v0::Constant>(
|
||||
ov::op::util::constant_predicate<int64_t>([](const std::vector<int64_t>& value) -> bool {
|
||||
return value.size() == 1 && (value[0] == 0 || value[0] == 1);
|
||||
}));
|
||||
auto gather_past_0 = wrap_type<ov::op::v8::Gather>({gather_input_0, beam_idx, axis_const});
|
||||
auto gather_past_1 = wrap_type<ov::op::v8::Gather>({gather_input_1, beam_idx, axis_const});
|
||||
auto kv_cache_0 = wrap_type<ov::intel_gpu::op::KVCache>({gather_past_0, any_input()});
|
||||
auto kv_cache_1 = wrap_type<ov::intel_gpu::op::KVCache>({gather_past_1, any_input()});
|
||||
|
||||
auto input_attn_mask = any_input();
|
||||
auto input_scale = any_input();
|
||||
auto sdpa_without_attn_mask_m = wrap_type<ov::intel_gpu::op::SDPA>({ any_input(), kv_cache_0, kv_cache_1 });
|
||||
auto sdpa_with_attn_mask_m = wrap_type<ov::intel_gpu::op::SDPA>({ any_input(), kv_cache_0, kv_cache_1, input_attn_mask });
|
||||
auto sdpa_with_attn_mask_and_scale_m =
|
||||
wrap_type<ov::intel_gpu::op::SDPA>({ any_input(), kv_cache_0, kv_cache_1, input_attn_mask, input_scale });
|
||||
|
||||
auto sdpa_m = std::make_shared<Or>(OutputVector{sdpa_without_attn_mask_m, sdpa_with_attn_mask_m, sdpa_with_attn_mask_and_scale_m});
|
||||
|
||||
ov::matcher_pass_callback callback = [OV_CAPTURE_CPY_AND_THIS](ov::pass::pattern::Matcher& m) {
|
||||
if (transformation_callback(m.get_match_root())) {
|
||||
return false;
|
||||
}
|
||||
const auto& pattern_map = m.get_pattern_value_map();
|
||||
|
||||
auto kv_cache_node_0 = std::dynamic_pointer_cast<ov::intel_gpu::op::KVCache>(pattern_map.at(kv_cache_0).get_node_shared_ptr());
|
||||
auto kv_cache_node_1 = std::dynamic_pointer_cast<ov::intel_gpu::op::KVCache>(pattern_map.at(kv_cache_1).get_node_shared_ptr());
|
||||
|
||||
auto beam_idx_node = pattern_map.at(beam_idx).get_node_shared_ptr();
|
||||
auto gather_input_node_0 = pattern_map.at(gather_input_0).get_node_shared_ptr();
|
||||
auto gather_input_node_1 = pattern_map.at(gather_input_1).get_node_shared_ptr();
|
||||
auto gather_node_0 = std::dynamic_pointer_cast<ov::op::v8::Gather>(pattern_map.at(gather_past_0).get_node_shared_ptr());
|
||||
auto gather_node_1 = std::dynamic_pointer_cast<ov::op::v8::Gather>(pattern_map.at(gather_past_1).get_node_shared_ptr());
|
||||
auto gather_axis_0 = gather_node_0->get_axis();
|
||||
auto gather_axis_1 = gather_node_1->get_axis();
|
||||
OPENVINO_ASSERT(gather_axis_0 == gather_axis_1);
|
||||
|
||||
ov::replace_node(gather_node_0, gather_input_node_0);
|
||||
ov::replace_node(gather_node_1, gather_input_node_1);
|
||||
|
||||
auto indirect_kv_cache_0 = std::make_shared<op::KVCache>(gather_input_node_0,
|
||||
kv_cache_node_0->get_input_node_shared_ptr(1),
|
||||
beam_idx_node,
|
||||
kv_cache_node_0->get_variable(),
|
||||
kv_cache_node_0->get_concat_axis(),
|
||||
gather_axis_0,
|
||||
kv_cache_node_0->get_output_element_type(0));
|
||||
|
||||
auto indirect_kv_cache_1 = std::make_shared<op::KVCache>(gather_input_node_1,
|
||||
kv_cache_node_1->get_input_node_shared_ptr(1),
|
||||
beam_idx_node,
|
||||
kv_cache_node_1->get_variable(),
|
||||
kv_cache_node_1->get_concat_axis(),
|
||||
gather_axis_1,
|
||||
kv_cache_node_1->get_output_element_type(0));
|
||||
|
||||
indirect_kv_cache_0->set_friendly_name(kv_cache_node_0->get_friendly_name());
|
||||
indirect_kv_cache_1->set_friendly_name(kv_cache_node_1->get_friendly_name());
|
||||
ov::copy_runtime_info(kv_cache_node_0, indirect_kv_cache_0);
|
||||
ov::copy_runtime_info(kv_cache_node_1, indirect_kv_cache_1);
|
||||
replace_node_unsafe(kv_cache_node_0, indirect_kv_cache_0);
|
||||
replace_node_unsafe(kv_cache_node_1, indirect_kv_cache_1);
|
||||
|
||||
auto sdpa = std::dynamic_pointer_cast<op::SDPA>(m.get_match_root());
|
||||
auto order_in0 = sdpa->get_input0_transpose_order();
|
||||
auto order_in1 = sdpa->get_input1_transpose_order();
|
||||
auto order_in2 = sdpa->get_input2_transpose_order();
|
||||
auto order_out = sdpa->get_output_transpose_order();
|
||||
auto is_causal = sdpa->get_causal();
|
||||
|
||||
std::shared_ptr<op::SDPA> indirect_sdpa;
|
||||
if (pattern_map.find(sdpa_without_attn_mask_m) != pattern_map.end()) {
|
||||
indirect_sdpa = std::make_shared<ov::intel_gpu::op::IndirectSDPA>(sdpa->get_input_node_shared_ptr(0),
|
||||
sdpa->get_input_node_shared_ptr(1),
|
||||
sdpa->get_input_node_shared_ptr(2),
|
||||
indirect_kv_cache_0->output(1), // beam table
|
||||
is_causal,
|
||||
gather_axis_1,
|
||||
order_in0,
|
||||
order_in1,
|
||||
order_in2,
|
||||
order_out);
|
||||
} else if (pattern_map.find(sdpa_with_attn_mask_m) != pattern_map.end()) {
|
||||
indirect_sdpa = std::make_shared<ov::intel_gpu::op::IndirectSDPA>(sdpa->get_input_node_shared_ptr(0),
|
||||
sdpa->get_input_node_shared_ptr(1),
|
||||
sdpa->get_input_node_shared_ptr(2),
|
||||
sdpa->get_input_node_shared_ptr(3),
|
||||
indirect_kv_cache_0->output(1), // beam table
|
||||
is_causal,
|
||||
gather_axis_1,
|
||||
order_in0,
|
||||
order_in1,
|
||||
order_in2,
|
||||
order_out);
|
||||
} else if (pattern_map.find(sdpa_with_attn_mask_and_scale_m) != pattern_map.end()) {
|
||||
indirect_sdpa = std::make_shared<ov::intel_gpu::op::IndirectSDPA>(sdpa->get_input_node_shared_ptr(0),
|
||||
sdpa->get_input_node_shared_ptr(1),
|
||||
sdpa->get_input_node_shared_ptr(2),
|
||||
sdpa->get_input_node_shared_ptr(3),
|
||||
sdpa->get_input_node_shared_ptr(4),
|
||||
indirect_kv_cache_0->output(1), // beam table
|
||||
is_causal,
|
||||
gather_axis_1,
|
||||
order_in0,
|
||||
order_in1,
|
||||
order_in2,
|
||||
order_out);
|
||||
}
|
||||
|
||||
OPENVINO_ASSERT(indirect_sdpa != nullptr);
|
||||
|
||||
indirect_sdpa->set_friendly_name(sdpa->get_friendly_name());
|
||||
ov::copy_runtime_info(sdpa, indirect_sdpa);
|
||||
ov::replace_node(sdpa, indirect_sdpa);
|
||||
|
||||
return true;
|
||||
};
|
||||
|
||||
auto m = std::make_shared<ov::pass::pattern::Matcher>(sdpa_m, "IndirectSDPAOpt");
|
||||
this->register_matcher(m, callback);
|
||||
}
|
||||
|
||||
IndirectKVCache::IndirectKVCache() {
|
||||
add_matcher<IndirectGemmOpt>();
|
||||
add_matcher<IndirectSDPAOpt>();
|
||||
}
|
||||
} // namespace intel_gpu
|
||||
} // namespace ov
|
||||
|
|
|
|||
|
|
@ -36,11 +36,22 @@ namespace intel_gpu {
|
|||
/// ┌────┴──────┐ ┌────┴──────┴───┐
|
||||
/// │ Gemm │ | IndirectGemm |
|
||||
/// └───────────┘ └───────────────┘
|
||||
class IndirectKVCache : public ov::pass::MatcherPass {
|
||||
class IndirectKVCache : public ov::pass::GraphRewrite {
|
||||
public:
|
||||
OPENVINO_RTTI("IndirectKVCache", "0");
|
||||
IndirectKVCache();
|
||||
};
|
||||
|
||||
class IndirectGemmOpt : public ov::pass::MatcherPass {
|
||||
public:
|
||||
OPENVINO_RTTI("IndirectGemmOpt", "0");
|
||||
IndirectGemmOpt();
|
||||
};
|
||||
|
||||
class IndirectSDPAOpt : public ov::pass::MatcherPass {
|
||||
public:
|
||||
OPENVINO_RTTI("IndirectSDPAOpt", "0");
|
||||
IndirectSDPAOpt();
|
||||
};
|
||||
} // namespace intel_gpu
|
||||
} // namespace ov
|
||||
|
|
|
|||
|
|
@ -0,0 +1,113 @@
|
|||
// Copyright (C) 2024 Intel Corporation
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
//
|
||||
|
||||
#include "intel_gpu/op/indirect_sdpa.hpp"
|
||||
#include "openvino/core/partial_shape.hpp"
|
||||
|
||||
namespace ov {
|
||||
namespace intel_gpu {
|
||||
namespace op {
|
||||
|
||||
IndirectSDPA::IndirectSDPA(const ov::Output<Node>& Q,
|
||||
const ov::Output<Node>& K,
|
||||
const ov::Output<Node>& V,
|
||||
const ov::Output<Node>& beam_table,
|
||||
const bool is_causal,
|
||||
const int64_t indirect_axis,
|
||||
const std::vector<int64_t>& order_q,
|
||||
const std::vector<int64_t>& order_k,
|
||||
const std::vector<int64_t>& order_v,
|
||||
const std::vector<int64_t>& order_out,
|
||||
const ov::element::Type output_type)
|
||||
: ov::intel_gpu::op::SDPA(Q, K, V, order_q, order_k, order_v, order_out, is_causal, output_type)
|
||||
, m_indirect_axis(indirect_axis) {
|
||||
set_argument(3, beam_table);
|
||||
validate_and_infer_types();
|
||||
}
|
||||
|
||||
IndirectSDPA::IndirectSDPA(const ov::Output<Node>& Q,
|
||||
const ov::Output<Node>& K,
|
||||
const ov::Output<Node>& V,
|
||||
const ov::Output<Node>& attn_mask,
|
||||
const ov::Output<Node>& beam_table,
|
||||
const bool is_causal,
|
||||
const int64_t indirect_axis,
|
||||
const std::vector<int64_t>& order_q,
|
||||
const std::vector<int64_t>& order_k,
|
||||
const std::vector<int64_t>& order_v,
|
||||
const std::vector<int64_t>& order_out,
|
||||
const ov::element::Type output_type)
|
||||
: ov::intel_gpu::op::SDPA(Q, K, V, attn_mask, order_q, order_k, order_v, order_out, is_causal, output_type)
|
||||
, m_indirect_axis(indirect_axis) {
|
||||
set_argument(4, beam_table);
|
||||
validate_and_infer_types();
|
||||
}
|
||||
|
||||
IndirectSDPA::IndirectSDPA(const ov::Output<Node>& Q,
|
||||
const ov::Output<Node>& K,
|
||||
const ov::Output<Node>& V,
|
||||
const ov::Output<Node>& attn_mask,
|
||||
const ov::Output<Node>& scale,
|
||||
const ov::Output<Node>& beam_table,
|
||||
const bool is_causal,
|
||||
const int64_t indirect_axis,
|
||||
const std::vector<int64_t>& order_q,
|
||||
const std::vector<int64_t>& order_k,
|
||||
const std::vector<int64_t>& order_v,
|
||||
const std::vector<int64_t>& order_out,
|
||||
const ov::element::Type output_type)
|
||||
: ov::intel_gpu::op::SDPA(Q, K, V, attn_mask, scale, order_q, order_k, order_v, order_out, is_causal, output_type)
|
||||
, m_indirect_axis(indirect_axis) {
|
||||
set_argument(5, beam_table);
|
||||
validate_and_infer_types();
|
||||
}
|
||||
|
||||
std::shared_ptr<ov::Node> IndirectSDPA::clone_with_new_inputs(const ov::OutputVector& new_args) const {
|
||||
check_new_args_count(this, new_args);
|
||||
|
||||
if (new_args.size() == 4) {
|
||||
return std::make_shared<IndirectSDPA>(new_args.at(0), new_args.at(1), new_args.at(2), new_args.at(3),
|
||||
m_is_causal, m_indirect_axis, m_order_q, m_order_k, m_order_v, m_order_out, m_output_type);
|
||||
} else if (new_args.size() == 5) {
|
||||
return std::make_shared<IndirectSDPA>(new_args.at(0), new_args.at(1), new_args.at(2), new_args.at(3), new_args.at(4),
|
||||
m_is_causal, m_indirect_axis, m_order_q, m_order_k, m_order_v, m_order_out, m_output_type);
|
||||
} else {
|
||||
return std::make_shared<IndirectSDPA>(new_args.at(0), new_args.at(1), new_args.at(2), new_args.at(3), new_args.at(4), new_args.at(5),
|
||||
m_is_causal, m_indirect_axis, m_order_q, m_order_k, m_order_v, m_order_out, m_output_type);
|
||||
}
|
||||
}
|
||||
|
||||
void IndirectSDPA::validate_and_infer_types() {
|
||||
const auto input_size = get_input_size();
|
||||
NODE_VALIDATION_CHECK(this,
|
||||
input_size == 4 || input_size == 5 || input_size == 6,
|
||||
"Number of inputs is incorrect. Current value is: ",
|
||||
input_size,
|
||||
", expected 4, 5 or 6.");
|
||||
|
||||
std::vector<ov::PartialShape> input_shapes;
|
||||
for (size_t i = 0; i < input_size - 1; i++) {
|
||||
input_shapes.push_back(get_input_partial_shape(i));
|
||||
}
|
||||
|
||||
auto out_shapes = shape_infer(this,
|
||||
input_shapes,
|
||||
m_order_q,
|
||||
m_order_k,
|
||||
m_order_v,
|
||||
m_order_out);
|
||||
|
||||
auto output_type = m_output_type == ov::element::undefined ? get_input_element_type(0) : m_output_type;
|
||||
set_output_type(0, output_type, out_shapes[0]);
|
||||
}
|
||||
|
||||
bool IndirectSDPA::visit_attributes(ov::AttributeVisitor &visitor) {
|
||||
SDPA::visit_attributes(visitor);
|
||||
visitor.on_attribute("indirect_axis", m_indirect_axis);
|
||||
return true;
|
||||
}
|
||||
|
||||
} // namespace op
|
||||
} // namespace intel_gpu
|
||||
} // namespace ov
|
||||
|
|
@ -30,6 +30,7 @@ SDPA::SDPA(const ov::Output<Node>& Q,
|
|||
, m_is_causal(is_causal)
|
||||
, m_output_type(output_type) {
|
||||
set_arguments({Q, K, V});
|
||||
set_causal(is_causal);
|
||||
validate_and_infer_types();
|
||||
}
|
||||
|
||||
|
|
@ -50,6 +51,7 @@ SDPA::SDPA(const ov::Output<Node>& Q,
|
|||
, m_is_causal(is_causal)
|
||||
, m_output_type(output_type) {
|
||||
set_arguments({Q, K, V, attn_mask});
|
||||
set_causal(is_causal);
|
||||
validate_and_infer_types();
|
||||
}
|
||||
|
||||
|
|
@ -71,13 +73,23 @@ SDPA::SDPA(const ov::Output<Node>& Q,
|
|||
, m_is_causal(is_causal)
|
||||
, m_output_type(output_type) {
|
||||
set_arguments({Q, K, V, attn_mask, scale});
|
||||
set_causal(is_causal);
|
||||
validate_and_infer_types();
|
||||
}
|
||||
|
||||
std::shared_ptr<ov::Node> SDPA::clone_with_new_inputs(const ov::OutputVector& new_args) const {
|
||||
check_new_args_count(this, new_args);
|
||||
|
||||
return std::make_shared<SDPA>(new_args.at(0), new_args.at(1), new_args.at(2), m_order_q, m_order_k, m_order_v, m_order_out, m_is_causal, m_output_type);
|
||||
if (new_args.size() == 3) {
|
||||
return std::make_shared<SDPA>(new_args.at(0), new_args.at(1), new_args.at(2),
|
||||
m_order_q, m_order_k, m_order_v, m_order_out, m_is_causal, m_output_type);
|
||||
} else if (new_args.size() == 4) {
|
||||
return std::make_shared<SDPA>(new_args.at(0), new_args.at(1), new_args.at(2), new_args.at(3),
|
||||
m_order_q, m_order_k, m_order_v, m_order_out, m_is_causal, m_output_type);
|
||||
} else {
|
||||
return std::make_shared<SDPA>(new_args.at(0), new_args.at(1), new_args.at(2), new_args.at(3), new_args.at(4),
|
||||
m_order_q, m_order_k, m_order_v, m_order_out, m_is_causal, m_output_type);
|
||||
}
|
||||
}
|
||||
|
||||
void SDPA::validate_and_infer_types() {
|
||||
|
|
|
|||
|
|
@ -80,14 +80,7 @@ TransposeSDPAMatcher::TransposeSDPAMatcher() {
|
|||
ov::matcher_pass_callback callback = [OV_CAPTURE_CPY_AND_THIS](Matcher& m) {
|
||||
const auto& pattern_map = m.get_pattern_value_map();
|
||||
|
||||
std::shared_ptr<ov::op::v13::ScaledDotProductAttention> sdpa;
|
||||
if (pattern_map.find(sdpa_without_attn_mask_m) != pattern_map.end()) {
|
||||
sdpa = std::dynamic_pointer_cast<ov::op::v13::ScaledDotProductAttention>(pattern_map.at(sdpa_without_attn_mask_m).get_node_shared_ptr());
|
||||
} else if (pattern_map.find(sdpa_with_attn_mask_m) != pattern_map.end()) {
|
||||
sdpa = std::dynamic_pointer_cast<ov::op::v13::ScaledDotProductAttention>(pattern_map.at(sdpa_with_attn_mask_m).get_node_shared_ptr());
|
||||
} else if (pattern_map.find(sdpa_with_attn_mask_and_scale_m) != pattern_map.end()) {
|
||||
sdpa = std::dynamic_pointer_cast<ov::op::v13::ScaledDotProductAttention>(pattern_map.at(sdpa_with_attn_mask_and_scale_m).get_node_shared_ptr());
|
||||
}
|
||||
auto sdpa = std::dynamic_pointer_cast<ov::op::v13::ScaledDotProductAttention>(m.get_match_root());
|
||||
|
||||
if (!sdpa || transformation_callback(sdpa)) {
|
||||
return false;
|
||||
|
|
@ -101,33 +94,41 @@ TransposeSDPAMatcher::TransposeSDPAMatcher() {
|
|||
size_t input_k_output_idx = sdpa->get_input_source_output(1).get_index();
|
||||
size_t input_v_output_idx = sdpa->get_input_source_output(2).get_index();
|
||||
|
||||
if (pattern_map.count(transpose_q_m) > 0) {
|
||||
auto tranpose_a_order = std::dynamic_pointer_cast<ov::op::v0::Constant>(pattern_map.at(transpose_q_order_m).get_node_shared_ptr());
|
||||
order_q = tranpose_a_order->cast_vector<int64_t>();
|
||||
if (order_q.back() != static_cast<int64_t>(order_q.size() - 1)) // Allow any transposes without head_size dim position change
|
||||
auto process_transpose = [](const std::shared_ptr<Node>& transpose_node,
|
||||
const std::shared_ptr<Node>& transpose_order_const_node,
|
||||
std::vector<int64_t>& order,
|
||||
size_t& output_idx) {
|
||||
auto transpose_order_const = std::dynamic_pointer_cast<ov::op::v0::Constant>(transpose_order_const_node);
|
||||
|
||||
order = transpose_order_const->cast_vector<int64_t>();
|
||||
// Allow any transposes without head_size dim position change
|
||||
if (order.back() != static_cast<int64_t>(order.size() - 1))
|
||||
return false;
|
||||
|
||||
auto tranpose_a = std::dynamic_pointer_cast<ov::op::v1::Transpose>(pattern_map.at(transpose_q_m).get_node_shared_ptr());
|
||||
input_q_output_idx = tranpose_a->get_input_source_output(0).get_index();
|
||||
}
|
||||
if (pattern_map.count(transpose_k_m) > 0) {
|
||||
auto tranpose_b_order = std::dynamic_pointer_cast<ov::op::v0::Constant>(pattern_map.at(transpose_k_order_m).get_node_shared_ptr());
|
||||
order_k = tranpose_b_order->cast_vector<int64_t>();
|
||||
if (order_k.back() != static_cast<int64_t>(order_k.size() - 1)) // Allow any transposes without head_size dim position change
|
||||
return false;
|
||||
auto transpose = std::dynamic_pointer_cast<ov::op::v1::Transpose>(transpose_node);
|
||||
output_idx = transpose->get_input_source_output(0).get_index();
|
||||
|
||||
auto tranpose_b = std::dynamic_pointer_cast<ov::op::v1::Transpose>(pattern_map.at(transpose_k_m).get_node_shared_ptr());
|
||||
input_k_output_idx = tranpose_b->get_input_source_output(0).get_index();
|
||||
}
|
||||
if (pattern_map.count(transpose_v_m) > 0) {
|
||||
auto tranpose_c_order = std::dynamic_pointer_cast<ov::op::v0::Constant>(pattern_map.at(transpose_v_order_m).get_node_shared_ptr());
|
||||
order_v = tranpose_c_order->cast_vector<int64_t>();
|
||||
if (order_v.back() != static_cast<int64_t>(order_v.size() - 1)) // Allow any transposes without head_size dim position change
|
||||
return false;
|
||||
return true;
|
||||
};
|
||||
|
||||
auto tranpose_c = std::dynamic_pointer_cast<ov::op::v1::Transpose>(pattern_map.at(transpose_k_m).get_node_shared_ptr());
|
||||
input_v_output_idx = tranpose_c->get_input_source_output(0).get_index();
|
||||
}
|
||||
bool can_fuse_transposes = true;
|
||||
if (pattern_map.count(transpose_q_m) > 0)
|
||||
can_fuse_transposes &= process_transpose(pattern_map.at(transpose_q_m).get_node_shared_ptr(),
|
||||
pattern_map.at(transpose_q_order_m).get_node_shared_ptr(),
|
||||
order_q, input_q_output_idx);
|
||||
|
||||
if (pattern_map.count(transpose_k_m) > 0)
|
||||
can_fuse_transposes &= process_transpose(pattern_map.at(transpose_k_m).get_node_shared_ptr(),
|
||||
pattern_map.at(transpose_k_order_m).get_node_shared_ptr(),
|
||||
order_k, input_k_output_idx);
|
||||
|
||||
if (pattern_map.count(transpose_v_m) > 0)
|
||||
can_fuse_transposes &= process_transpose(pattern_map.at(transpose_v_m).get_node_shared_ptr(),
|
||||
pattern_map.at(transpose_v_order_m).get_node_shared_ptr(),
|
||||
order_v, input_v_output_idx);
|
||||
|
||||
if (!can_fuse_transposes)
|
||||
return false;
|
||||
|
||||
auto input_q = ov::Output<Node>(pattern_map.at(input_q_m).get_node_shared_ptr(), input_q_output_idx);
|
||||
auto input_k = ov::Output<Node>(pattern_map.at(input_k_m).get_node_shared_ptr(), input_k_output_idx);
|
||||
|
|
|
|||
|
|
@ -13,6 +13,7 @@
|
|||
#include <vector>
|
||||
|
||||
#include "intel_gpu/plugin/transformations_pipeline.hpp"
|
||||
#include "intel_gpu/runtime/debug_configuration.hpp"
|
||||
#include "intel_gpu/runtime/itt.hpp"
|
||||
#include "low_precision/convolution.hpp"
|
||||
#include "low_precision/convolution_backprop_data.hpp"
|
||||
|
|
@ -307,6 +308,10 @@ void TransformationsPipeline::apply(std::shared_ptr<ov::Model> func) {
|
|||
manager.register_pass<ov::pass::CommonOptimizations>();
|
||||
|
||||
pass_config->set_callback<ov::pass::ScaledDotProductAttentionDecomposition>([&](const std::shared_ptr<const ov::Node> node){
|
||||
GPU_DEBUG_IF(cldnn::debug_configuration::get_instance()->enable_sdpa != -1) {
|
||||
GPU_DEBUG_CODE(return cldnn::debug_configuration::get_instance()->enable_sdpa == 1);
|
||||
}
|
||||
|
||||
if (!config.get_property(ov::intel_gpu::hint::enable_sdpa_optimization))
|
||||
return false;
|
||||
|
||||
|
|
|
|||
|
|
@ -171,6 +171,8 @@ static void print_help_messages() {
|
|||
message_list.emplace_back("OV_GPU_DisableDynamicImpl", "Disable dynamic implementation");
|
||||
message_list.emplace_back("OV_GPU_DisableRuntimeBufferFusing", "Disable runtime buffer fusing");
|
||||
message_list.emplace_back("OV_GPU_DisableMemoryReuse", "Disable memory reuse");
|
||||
message_list.emplace_back("OV_GPU_EnableSDPA", "This allows the enforcement of SDPA decomposition logic: 0 completely disables SDPA kernel usage, "
|
||||
"and 1 enables it for all the cases.");
|
||||
message_list.emplace_back("OV_GPU_DumpMemoryPool", "Dump memory pool contents of each iteration");
|
||||
message_list.emplace_back("OV_GPU_DumpMemoryPoolIters", "List of iterations to dump memory pool status, separated by space.");
|
||||
message_list.emplace_back("OV_GPU_DumpMemoryPoolPath", "Enable dumping memory pool status to csv file and set the dest path");
|
||||
|
|
@ -232,6 +234,7 @@ debug_configuration::debug_configuration()
|
|||
, serialize_compile(0)
|
||||
, max_kernels_per_batch(0)
|
||||
, impls_cache_capacity(-1)
|
||||
, enable_sdpa(-1)
|
||||
, disable_async_compilation(0)
|
||||
, disable_winograd_conv(0)
|
||||
, disable_dynamic_impl(0)
|
||||
|
|
@ -280,6 +283,7 @@ debug_configuration::debug_configuration()
|
|||
get_gpu_debug_env_var("ForceImplTypes", forced_impl_types_str);
|
||||
get_gpu_debug_env_var("MaxKernelsPerBatch", max_kernels_per_batch);
|
||||
get_gpu_debug_env_var("ImplsCacheCapacity", impls_cache_capacity);
|
||||
get_gpu_debug_env_var("EnableSDPA", enable_sdpa);
|
||||
get_gpu_debug_env_var("DisableAsyncCompilation", disable_async_compilation);
|
||||
get_gpu_debug_env_var("DisableWinogradConv", disable_winograd_conv);
|
||||
get_gpu_debug_env_var("DisableDynamicImpl", disable_dynamic_impl);
|
||||
|
|
|
|||
|
|
@ -0,0 +1,40 @@
|
|||
// Copyright (C) 2018-2024 Intel Corporation
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
//
|
||||
#include "single_op_tests/roi_align_rotated.hpp"
|
||||
|
||||
#include "common_test_utils/test_constants.hpp"
|
||||
|
||||
namespace {
|
||||
using ov::test::ROIAlignRotatedLayerTest;
|
||||
|
||||
const std::vector<ov::element::Type> netPRCs = {
|
||||
ov::element::f32
|
||||
// There is no possibility to test ROIAlign in fp16 precision,
|
||||
// because on edge cases where in fp32 version ROI value is
|
||||
// a little bit smaller than the nearest integer value,
|
||||
// it would be bigger than the nearest integer in fp16 precision.
|
||||
// Such behavior leads to completely different results of ROIAlign
|
||||
// in fp32 and fp16 precisions.
|
||||
// In real AI applications this problem is solved by precision-aware training.
|
||||
|
||||
// ov::element::f16
|
||||
};
|
||||
|
||||
INSTANTIATE_TEST_SUITE_P(gtest_smoke_TestsROIAlignRotatedROIAlignLayerTest_EvalGenerateName_,
|
||||
ROIAlignRotatedLayerTest,
|
||||
::testing::Combine(::testing::ValuesIn(ov::test::static_shapes_to_test_representation(
|
||||
std::vector<std::vector<ov::Shape>>{{{3, 8, 16, 16}},
|
||||
{{2, 1, 16, 10}},
|
||||
{{4, 3, 5, 12}}})),
|
||||
::testing::ValuesIn(std::vector<int>{2, 4}),
|
||||
::testing::Values(2),
|
||||
::testing::Values(2),
|
||||
::testing::Values(2),
|
||||
::testing::ValuesIn(std::vector<float>{1, 0.625}),
|
||||
::testing::ValuesIn(std::vector<bool>{true, false}),
|
||||
::testing::ValuesIn(netPRCs),
|
||||
::testing::Values(ov::test::utils::DEVICE_GPU)),
|
||||
ROIAlignRotatedLayerTest::getTestCaseName);
|
||||
|
||||
} // namespace
|
||||
|
|
@ -420,6 +420,7 @@ std::vector<ov::ProfilingInfo> ZeroInferRequest::get_profiling_info() const {
|
|||
const auto& compiledModel = *std::dynamic_pointer_cast<const ICompiledModel>(_compiledModel);
|
||||
const auto& compilerConfig = compiledModel.get_config();
|
||||
if (!compilerConfig.get<PERF_COUNT>() || !_config.get<PERF_COUNT>()) {
|
||||
_logger.debug("InferRequest::get_profiling_info complete with empty {}.");
|
||||
return {};
|
||||
}
|
||||
|
||||
|
|
@ -431,16 +432,18 @@ std::vector<ov::ProfilingInfo> ZeroInferRequest::get_profiling_info() const {
|
|||
const auto& compiler = compiledModel.get_compiler();
|
||||
const auto& blob = networkDesc->compiledNetwork;
|
||||
auto profData = get_raw_profiling_data();
|
||||
_logger.debug("InferRequest::get_profiling_info complete with compiler->process_profiling_output().");
|
||||
return compiler->process_profiling_output(profData, blob, compilerConfig);
|
||||
} else {
|
||||
auto proftype = _config.get<PROFILING_TYPE>();
|
||||
if (proftype == ov::intel_npu::ProfilingType::INFER) {
|
||||
_logger.debug("InferRequest::get_profiling_info complete with _npuProfiling->getNpuInferStatistics().");
|
||||
return _npuProfiling->getNpuInferStatistics();
|
||||
} else { /// proftype = MODEL or undefined = fallback to model profiling
|
||||
_logger.debug("InferRequest::get_profiling_info complete with _profilingQuery.getLayerStatistics().");
|
||||
return _profilingQuery.getLayerStatistics();
|
||||
}
|
||||
}
|
||||
_logger.debug("InferRequest::get_profiling_info completed");
|
||||
}
|
||||
|
||||
std::vector<uint8_t> ZeroInferRequest::get_raw_profiling_data() const {
|
||||
|
|
|
|||
|
|
@ -106,7 +106,7 @@ private:
|
|||
const ov::element::Type precision = ov::element::f32;
|
||||
|
||||
ov::ParameterVector params{std::make_shared<ov::op::v0::Parameter>(precision, ov::Shape{input_shape})};
|
||||
auto constant = ov::test::utils::deprecated::make_constant(precision, {4096, 1024}, std::vector<float>{}, true);
|
||||
auto constant = ov::test::utils::make_constant(precision, ov::Shape{4096, 1024});
|
||||
auto custom_op = std::make_shared<UnsupportedTestOperation>(constant);
|
||||
|
||||
ov::NodeVector results{custom_op};
|
||||
|
|
|
|||
|
|
@ -51,7 +51,7 @@ std::shared_ptr<ov::Model> createModelWithUnknownNode() {
|
|||
const ov::Shape input_shape = {1, 4096};
|
||||
const ov::element::Type precision = ov::element::f32;
|
||||
ov::ParameterVector params = {std::make_shared<ov::op::v0::Parameter>(precision, ov::Shape{input_shape})};
|
||||
auto constant = ov::test::utils::deprecated::make_constant(precision, {4096, 1024}, std::vector<float>{}, true);
|
||||
auto constant = ov::test::utils::make_constant(precision, ov::Shape{4096, 1024});
|
||||
auto custom_op = std::make_shared<UnsupportedTestOp>(constant);
|
||||
|
||||
ov::NodeVector results{custom_op};
|
||||
|
|
|
|||
|
|
@ -0,0 +1,15 @@
|
|||
// Copyright (C) 2018-2024 Intel Corporation
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
//
|
||||
|
||||
#pragma once
|
||||
|
||||
#include "shared_test_classes/single_op/roi_align_rotated.hpp"
|
||||
|
||||
namespace ov {
|
||||
namespace test {
|
||||
TEST_P(ROIAlignRotatedLayerTest, Inference) {
|
||||
run();
|
||||
}
|
||||
} // namespace test
|
||||
} // namespace ov
|
||||
|
|
@ -0,0 +1,30 @@
|
|||
// Copyright (C) 2018-2024 Intel Corporation
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
//
|
||||
|
||||
#pragma once
|
||||
|
||||
#include "shared_test_classes/base/ov_subgraph.hpp"
|
||||
|
||||
namespace ov {
|
||||
namespace test {
|
||||
using roialignrotatedParams = std::tuple<std::vector<InputShape>, // Feature map shape
|
||||
int, // Num of Rois
|
||||
int, // Pooled h
|
||||
int, // Pooled w
|
||||
int, // Sampling ratio
|
||||
float, // Spatial scale
|
||||
bool, // Clockwise mode
|
||||
ov::element::Type, // Model type
|
||||
ov::test::TargetDevice>; // Device name
|
||||
|
||||
class ROIAlignRotatedLayerTest : public testing::WithParamInterface<roialignrotatedParams>,
|
||||
virtual public ov::test::SubgraphBaseTest {
|
||||
public:
|
||||
static std::string getTestCaseName(const testing::TestParamInfo<roialignrotatedParams>& obj);
|
||||
|
||||
protected:
|
||||
void SetUp() override;
|
||||
};
|
||||
} // namespace test
|
||||
} // namespace ov
|
||||
|
|
@ -0,0 +1,138 @@
|
|||
// Copyright (C) 2018-2024 Intel Corporation
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
//
|
||||
|
||||
#include "shared_test_classes/single_op/roi_align_rotated.hpp"
|
||||
|
||||
#include <random>
|
||||
|
||||
#include "openvino/core/enum_names.hpp"
|
||||
|
||||
namespace ov {
|
||||
namespace test {
|
||||
|
||||
static constexpr int ROI_DEF_SIZE = 5;
|
||||
static constexpr int SEED = 7877;
|
||||
static constexpr float PI = 3.14159265358979323846f;
|
||||
|
||||
struct TestParams {
|
||||
std::vector<InputShape> input_shapes;
|
||||
int num_rois;
|
||||
int pooled_h;
|
||||
int pooled_w;
|
||||
int sampliong_ratio;
|
||||
float spatial_scale;
|
||||
bool clockwise_mode;
|
||||
ov::element::Type model_type;
|
||||
std::string target_device;
|
||||
};
|
||||
|
||||
static TestParams ExtractTestParams(const roialignrotatedParams& param) {
|
||||
TestParams tp;
|
||||
std::tie(tp.input_shapes,
|
||||
tp.num_rois,
|
||||
tp.pooled_h,
|
||||
tp.pooled_w,
|
||||
tp.sampliong_ratio,
|
||||
tp.spatial_scale,
|
||||
tp.clockwise_mode,
|
||||
tp.model_type,
|
||||
tp.target_device) = param;
|
||||
return tp;
|
||||
}
|
||||
|
||||
static float RandomFloat(float low, float high) {
|
||||
static std::default_random_engine engine(SEED);
|
||||
std::uniform_real_distribution<float> dis(low, high);
|
||||
return dis(engine);
|
||||
}
|
||||
|
||||
static std::vector<float> FillRoisTensor(int num_rois, int height, int width) {
|
||||
std::vector<float> rois;
|
||||
rois.resize(num_rois * ROI_DEF_SIZE);
|
||||
|
||||
for (int i = 0; i < rois.size() / ROI_DEF_SIZE; i++) {
|
||||
// center_x, center_y, width, height, angle
|
||||
rois[i * ROI_DEF_SIZE + 0] = RandomFloat(0.0f, width);
|
||||
rois[i * ROI_DEF_SIZE + 1] = RandomFloat(0.0f, height);
|
||||
rois[i * ROI_DEF_SIZE + 2] = RandomFloat(0.0f, width);
|
||||
rois[i * ROI_DEF_SIZE + 3] = RandomFloat(0.0f, height);
|
||||
rois[i * ROI_DEF_SIZE + 4] = RandomFloat(0.0f, 2 * PI);
|
||||
}
|
||||
|
||||
return rois;
|
||||
}
|
||||
|
||||
static std::vector<int> FillBAtchIdxTensor(int num_rois, int batch_size) {
|
||||
std::vector<int> idx;
|
||||
idx.resize(num_rois);
|
||||
int batch_id = 0;
|
||||
for (int i = 0; i < idx.size(); i++) {
|
||||
idx[i] = batch_id;
|
||||
batch_id = (batch_id + 1) % batch_size;
|
||||
}
|
||||
|
||||
return idx;
|
||||
}
|
||||
|
||||
std::string ROIAlignRotatedLayerTest::getTestCaseName(const testing::TestParamInfo<roialignrotatedParams>& obj) {
|
||||
const TestParams tp = ExtractTestParams(obj.param);
|
||||
|
||||
std::ostringstream result;
|
||||
result << "IS=(";
|
||||
for (size_t i = 0lu; i < tp.input_shapes.size(); i++) {
|
||||
result << ov::test::utils::partialShape2str({tp.input_shapes[i].first})
|
||||
<< (i < tp.input_shapes.size() - 1lu ? "_" : "");
|
||||
}
|
||||
result << ")_TS=";
|
||||
for (size_t i = 0lu; i < tp.input_shapes.front().second.size(); i++) {
|
||||
result << "{";
|
||||
for (size_t j = 0lu; j < tp.input_shapes.size(); j++) {
|
||||
result << ov::test::utils::vec2str(tp.input_shapes[j].second[i])
|
||||
<< (j < tp.input_shapes.size() - 1lu ? "_" : "");
|
||||
}
|
||||
result << "}_";
|
||||
}
|
||||
result << "numRois=" << tp.num_rois << "_";
|
||||
result << "pooledH=" << tp.pooled_h << "_";
|
||||
result << "pooledW=" << tp.pooled_w << "_";
|
||||
result << "samplingRatio=" << tp.sampliong_ratio << "_";
|
||||
result << "spatialScale=" << tp.spatial_scale << "_";
|
||||
result << "clockwiseMode=" << tp.clockwise_mode << "_";
|
||||
result << "modelType=" << tp.model_type.to_string() << "_";
|
||||
result << "trgDev=" << tp.target_device;
|
||||
return result.str();
|
||||
}
|
||||
|
||||
void ROIAlignRotatedLayerTest::SetUp() {
|
||||
const TestParams tp = ExtractTestParams(this->GetParam());
|
||||
targetDevice = tp.target_device;
|
||||
init_input_shapes(tp.input_shapes);
|
||||
|
||||
const auto input_batch_size = inputDynamicShapes[0][0].get_length();
|
||||
const auto input_height = inputDynamicShapes[0][2].get_length();
|
||||
const auto input_width = inputDynamicShapes[0][3].get_length();
|
||||
|
||||
auto input = std::make_shared<ov::op::v0::Parameter>(tp.model_type, inputDynamicShapes[0]);
|
||||
const auto rois_shape = ov::Shape{static_cast<size_t>(tp.num_rois), ROI_DEF_SIZE};
|
||||
const auto rois_idx_shape = ov::Shape{static_cast<size_t>(tp.num_rois)};
|
||||
|
||||
auto rois = std::make_shared<ov::op::v0::Constant>(tp.model_type,
|
||||
rois_shape,
|
||||
FillRoisTensor(tp.num_rois, input_height, input_width).data());
|
||||
auto rois_idx = std::make_shared<ov::op::v0::Constant>(ov::element::i32,
|
||||
rois_idx_shape,
|
||||
FillBAtchIdxTensor(tp.num_rois, input_batch_size).data());
|
||||
auto roi_align = std::make_shared<ov::op::v15::ROIAlignRotated>(input,
|
||||
rois,
|
||||
rois_idx,
|
||||
tp.pooled_h,
|
||||
tp.pooled_w,
|
||||
tp.sampliong_ratio,
|
||||
tp.spatial_scale,
|
||||
tp.clockwise_mode);
|
||||
function = std::make_shared<ov::Model>(roi_align->outputs(), ov::ParameterVector{input}, "roi_align_rotated");
|
||||
}
|
||||
|
||||
} // namespace test
|
||||
} // namespace ov
|
||||
|
|
@ -11,65 +11,6 @@
|
|||
namespace ov {
|
||||
namespace test {
|
||||
namespace utils {
|
||||
|
||||
namespace deprecated {
|
||||
// Legacy implementation
|
||||
// Remove after transition to new one
|
||||
template <typename T>
|
||||
std::shared_ptr<ov::Node> make_constant(const ov::element::Type& type,
|
||||
const std::vector<size_t>& shape,
|
||||
const std::vector<T>& data,
|
||||
bool random = false,
|
||||
T up_to = 10,
|
||||
T start_from = 1,
|
||||
const int seed = 1) {
|
||||
#define makeNode(TYPE) \
|
||||
case TYPE: \
|
||||
if (random) { \
|
||||
return std::make_shared<ov::op::v0::Constant>( \
|
||||
type, \
|
||||
shape, \
|
||||
generateVector<TYPE>(ov::shape_size(shape), \
|
||||
ov::element_type_traits<TYPE>::value_type(up_to), \
|
||||
ov::element_type_traits<TYPE>::value_type(start_from), \
|
||||
seed)); \
|
||||
} else { \
|
||||
if (std::is_same<T, fundamental_type_for<TYPE>>::value) { \
|
||||
return std::make_shared<ov::op::v0::Constant>(type, shape, data); \
|
||||
} else { \
|
||||
/* Convert std::vector<T> data to required type */ \
|
||||
std::vector<fundamental_type_for<TYPE>> converted_data(data.size()); \
|
||||
std::transform(data.cbegin(), data.cend(), converted_data.begin(), [](T e) { \
|
||||
return static_cast<fundamental_type_for<TYPE>>(e); \
|
||||
}); \
|
||||
return std::make_shared<ov::op::v0::Constant>(type, shape, converted_data); \
|
||||
} \
|
||||
} \
|
||||
break;
|
||||
switch (type) {
|
||||
makeNode(ov::element::bf16);
|
||||
makeNode(ov::element::f16);
|
||||
makeNode(ov::element::f32);
|
||||
makeNode(ov::element::f64);
|
||||
makeNode(ov::element::i8);
|
||||
makeNode(ov::element::i16);
|
||||
makeNode(ov::element::i32);
|
||||
makeNode(ov::element::i64);
|
||||
makeNode(ov::element::u8);
|
||||
makeNode(ov::element::u16);
|
||||
makeNode(ov::element::u32);
|
||||
makeNode(ov::element::u64);
|
||||
makeNode(ov::element::boolean);
|
||||
makeNode(ov::element::nf4);
|
||||
makeNode(ov::element::u4);
|
||||
makeNode(ov::element::i4);
|
||||
default:
|
||||
throw std::runtime_error("Unhandled precision");
|
||||
}
|
||||
#undef makeNode
|
||||
}
|
||||
} // namespace deprecated
|
||||
|
||||
std::shared_ptr<ov::Node> make_constant(const ov::element::Type& type,
|
||||
const ov::Shape& shape,
|
||||
InputGenerateData in_data = InputGenerateData(1, 9, 1000, 1));
|
||||
|
|
|
|||
|
|
@ -287,6 +287,28 @@ TEST_DATA(LIST(1, 1, 5, 5),
|
|||
LIST(0),
|
||||
LIST(5.1271, 1.2473, 6.1773, 2.9598, 7.2275, 3.2300, 8.2777, 3.7458, 9.3279, 4.4060),
|
||||
"roi_align_rotated_all_features");
|
||||
TEST_DATA(LIST(1, 1, 2, 5),
|
||||
2,
|
||||
2,
|
||||
1.0f,
|
||||
2,
|
||||
true,
|
||||
LIST(1, 2, 3, 4, 5, 6, 7, 8, 9, 10),
|
||||
LIST(0.5, 0.5, 1, 1, 0),
|
||||
LIST(0),
|
||||
LIST(1.0, 1.25, 2.25, 2.50),
|
||||
"input_image_not_rectangular");
|
||||
TEST_DATA(LIST(2, 1, 2, 5),
|
||||
2,
|
||||
2,
|
||||
1.0f,
|
||||
2,
|
||||
true,
|
||||
LIST(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20),
|
||||
LIST(0.5, 1., 2., 5., 0.5, 0., 2., 5., 1., 0.),
|
||||
LIST(0, 1),
|
||||
LIST(0.5201, 1.9866, 2.5219, 3.0896, 0.0000, 16.7500, 0.0000, 16.7500),
|
||||
"input_image_not_rectangular_batch_2");
|
||||
|
||||
#undef PI
|
||||
#undef LIST
|
||||
|
|
|
|||
|
|
@ -1016,6 +1016,28 @@ def create_pytorch_module_with_single_input_as_list(tmp_dir):
|
|||
"compress_to_fp16": False}
|
||||
|
||||
|
||||
def create_pytorch_module_with_nested_dict_input(tmp_dir):
|
||||
class PTModel(torch.nn.Module):
|
||||
def forward(self, a, b):
|
||||
return a["1"] * a["2"] + b
|
||||
|
||||
net = PTModel()
|
||||
a1 = ov.opset10.parameter(PartialShape([-1]), dtype=np.float32)
|
||||
a2 = ov.opset10.parameter(PartialShape([-1]), dtype=np.float32)
|
||||
b = ov.opset10.parameter(PartialShape([-1]), dtype=np.float32)
|
||||
mul = ov.opset10.multiply(a1, a2)
|
||||
add = ov.opset10.add(mul, b)
|
||||
ref_model = Model([add], [a1, a2, b], "test")
|
||||
return net, ref_model, {
|
||||
"example_input": (
|
||||
{
|
||||
"1": torch.tensor([1, 2], dtype=torch.float32),
|
||||
"2": torch.tensor([3, 4], dtype=torch.float32)
|
||||
},
|
||||
torch.tensor([5, 6], dtype=torch.float32)
|
||||
)}
|
||||
|
||||
|
||||
class TestMoConvertPyTorch(CommonMOConvertTest):
|
||||
test_data = [
|
||||
create_pytorch_nn_module_case1,
|
||||
|
|
@ -1067,7 +1089,8 @@ class TestMoConvertPyTorch(CommonMOConvertTest):
|
|||
create_pytorch_module_with_nested_inputs5,
|
||||
create_pytorch_module_with_nested_inputs6,
|
||||
create_pytorch_module_with_nested_list_and_single_input,
|
||||
create_pytorch_module_with_single_input_as_list
|
||||
create_pytorch_module_with_single_input_as_list,
|
||||
create_pytorch_module_with_nested_dict_input
|
||||
]
|
||||
|
||||
@pytest.mark.parametrize("create_model", test_data)
|
||||
|
|
|
|||
|
|
@ -157,7 +157,7 @@ class TestTFEqual(CommonTFLayerTest):
|
|||
pytest.param(
|
||||
dict(x_shape=[9], y_shape=[9], # Comparing shapes which contains corner cases
|
||||
x_value=x_corner, y_value=y_corner),
|
||||
marks=pytest.mark.xfail(reason="94234")),
|
||||
marks=pytest.mark.special_xfail(args={"ie_device": "GPU"}, reason="94234")),
|
||||
dict(x_shape=[1, 2, 3, 4], y_shape=[1, 2, 3, 4])
|
||||
# Comparing shapes with different dimensions (more than 3, for case with nchw/nhcw), random values (false and possible true)
|
||||
]
|
||||
|
|
@ -179,7 +179,7 @@ class TestTFEqual(CommonTFLayerTest):
|
|||
pytest.param(
|
||||
dict(x_shape=[9], y_shape=[9], # Comparing shapes which contains corner cases
|
||||
x_value=x_corner, y_value=y_corner),
|
||||
marks=pytest.mark.xfail(reason="94234")),
|
||||
marks=pytest.mark.special_xfail(args={"ie_device": "GPU"}, reason="94234")),
|
||||
dict(x_shape=[1, 2, 3, 4], y_shape=[1, 2, 3, 4])
|
||||
# Comparing shapes with different dimensions (more than 3, for case with nchw/nhcw), random values (false and possible true)
|
||||
]
|
||||
|
|
@ -201,7 +201,7 @@ class TestTFEqual(CommonTFLayerTest):
|
|||
pytest.param(
|
||||
dict(x_shape=[9], y_shape=[9], # Comparing shapes which contains corner cases
|
||||
x_value=x_corner, y_value=y_corner),
|
||||
marks=pytest.mark.xfail(reason="94234")),
|
||||
marks=pytest.mark.special_xfail(args={"ie_device": "GPU"}, reason="94234")),
|
||||
dict(x_shape=[1, 2, 3, 4], y_shape=[1, 2, 3, 4])
|
||||
# Comparing shapes with different dimensions (more than 3, for case with nchw/nhcw), random values (false and possible true)
|
||||
]
|
||||
|
|
|
|||
Loading…
Reference in New Issue