diff --git a/.github/scripts/workflow_rerun/errors_to_look_for.json b/.github/scripts/workflow_rerun/errors_to_look_for.json index b9cac8f17adaa6..d8fe6ac2df03d2 100644 --- a/.github/scripts/workflow_rerun/errors_to_look_for.json +++ b/.github/scripts/workflow_rerun/errors_to_look_for.json @@ -86,5 +86,25 @@ { "error_text": "because the GET request got Content-Type", "ticket": 158400 + }, + { + "error_text": "Unable to make request:", + "ticket": 158401 + }, + { + "error_text": "Failed to make request", + "ticket": 158401 + }, + { + "error_text": "Failure when receiving data from the peer", + "ticket": 159323 + }, + { + "error_text": "HTTP response code said error", + "ticket": 159398 + }, + { + "error_text": "download failed after attempts", + "ticket": 159547 } ] \ No newline at end of file diff --git a/docs/articles_en/about-openvino/compatibility-and-support/supported-models.rst b/docs/articles_en/about-openvino/compatibility-and-support/supported-models.rst index d877cb1768d44d..f4ec275491fa32 100644 --- a/docs/articles_en/about-openvino/compatibility-and-support/supported-models.rst +++ b/docs/articles_en/about-openvino/compatibility-and-support/supported-models.rst @@ -6,16 +6,14 @@ models from OpenVINO-supported frameworks may also work properly but have not be **AI Models that run on Intel® Core Ultra™ Processors with OpenVINO™ toolkit:** -.. raw:: html - - - - -.. csv-table:: +.. data-table:: :class: modeldata stripe :name: supportedModelsTable :header-rows: 1 :file: ../../_static/download/supported_models.csv + :data-column-hidden: [] + :data-order: [[ 0, "asc" ]] + :data-page-length: 10 | Marked cells indicate models that passed inference with no errors. Empty cells indicate diff --git a/docs/articles_en/about-openvino/compatibility-and-support/supported-operations.rst b/docs/articles_en/about-openvino/compatibility-and-support/supported-operations.rst index d27f7626391f46..1bd8f5dae7c634 100644 --- a/docs/articles_en/about-openvino/compatibility-and-support/supported-operations.rst +++ b/docs/articles_en/about-openvino/compatibility-and-support/supported-operations.rst @@ -41,27 +41,36 @@ Data as of OpenVINO 2024.4, 18 Oct. 2024. .. tab-item:: PyTorch - .. csv-table:: + .. data-table:: :class: modeldata stripe - :name: TensorFlow ops + :name: TensorFlow_ops_v1 :header-rows: 1 :file: ../../_static/conformance_files/pytorch_ops.csv + :data-column-hidden: [] + :data-order: [[ 0, "asc" ]] + :data-page-length: 10 .. tab-item:: TensorFlow - .. csv-table:: + .. data-table:: :class: modeldata stripe - :name: TensorFlow ops + :name: TensorFlow_ops_v2 :header-rows: 1 :file: ../../_static/conformance_files/tensorflow_ops.csv + :data-column-hidden: [] + :data-order: [[ 0, "asc" ]] + :data-page-length: 10 .. tab-item:: PaddlePaddle - .. csv-table:: + .. data-table:: :class: modeldata stripe - :name: Paddle ops + :name: Paddle_ops :header-rows: 1 :file: ../../_static/conformance_files/paddlepaddle_ops.csv + :data-column-hidden: [] + :data-order: [[ 0, "asc" ]] + :data-page-length: 10 .. tab-item:: ONNX diff --git a/docs/articles_en/about-openvino/performance-benchmarks/generative-ai-performance.rst b/docs/articles_en/about-openvino/performance-benchmarks/generative-ai-performance.rst index 085a1ff8449151..83581d465df92e 100644 --- a/docs/articles_en/about-openvino/performance-benchmarks/generative-ai-performance.rst +++ b/docs/articles_en/about-openvino/performance-benchmarks/generative-ai-performance.rst @@ -8,10 +8,6 @@ The current data is as of OpenVINO 2024.4, 20 Nov. 2024. The tables below list the key performance indicators for inference on built-in GPUs. -.. raw:: html - - - .. tab-set:: @@ -22,7 +18,9 @@ The tables below list the key performance indicators for inference on built-in G :name: supportedModelsTable_V1 :header-rows: 1 :file: ../../_static/benchmarks_files/llm_models_9-288V.csv - :hidden: [3,4,6] + :data-column-hidden: [3,4,6] + :data-order: [[ 0, "asc" ]] + :data-page-length: 10 .. tab-item:: 7-268V @@ -31,7 +29,8 @@ The tables below list the key performance indicators for inference on built-in G :name: supportedModelsTable_V2 :header-rows: 1 :file: ../../_static/benchmarks_files/llm_models_7-258V.csv - :hidden: [3,4,6] + :data-column-hidden: [3,4,6] + :data-order: [[ 0, "asc" ]] .. tab-item:: 7-155H @@ -40,7 +39,8 @@ The tables below list the key performance indicators for inference on built-in G :name: supportedModelsTable_V3 :header-rows: 1 :file: ../../_static/benchmarks_files/llm_models_7-155H.csv - :hidden: [3,4,6] + :data-column-hidden: [3,4,6] + :data-order: [[ 0, "asc" ]] .. grid:: 1 1 2 2 diff --git a/docs/openvino_sphinx_theme/openvino_sphinx_theme/directives/code.py b/docs/openvino_sphinx_theme/openvino_sphinx_theme/directives/code.py index c3e0e81eec3b3a..814517289ce114 100644 --- a/docs/openvino_sphinx_theme/openvino_sphinx_theme/directives/code.py +++ b/docs/openvino_sphinx_theme/openvino_sphinx_theme/directives/code.py @@ -11,7 +11,7 @@ import requests import re import json - +import html import csv logger = logging.getLogger(__name__) @@ -147,7 +147,9 @@ class DataTable(Directive): 'file': directives.path, 'class': directives.unchanged, 'name': directives.unchanged, - 'hidden': directives.unchanged + 'data-column-hidden': directives.unchanged, + 'data-page-length': directives.unchanged, + 'data-order': directives.unchanged } def run(self) -> List[Node]: @@ -159,10 +161,12 @@ def run(self) -> List[Node]: csv_node = [] with open(csv_file, 'r') as j: csv_data = list(csv.reader(j)) - class_table_tag = ' class="' + "".join(c for c in str(self.options['class']) + '"') if 'class' in self.options is not None else "" - id_table_tag = ' id="' + "".join(c for c in str(self.options['name']) + '"') if 'name' in self.options is not None else "" - hidden_table_tag = ' data-columns-hidden="' + "".join(c for c in str(self.options['hidden']) + '"') if 'hidden' in self.options is not None else "" - csv_table_html = '' + class_table_tag = f' class="{html.escape(self.options["class"])}"' if "class" in self.options else "" + id_table_tag = f' id="{html.escape(self.options["name"])}"' if "name" in self.options else "" + data_column_hidden_tag = f' data-column-hidden="{html.escape(self.options["data-column-hidden"])}"' if "data-column-hidden" in self.options else "" + data_order_tag = f' data-order="{html.escape(self.options["data-order"])}"' if "data-order" in self.options else "" + data_page_length_tag = f' data-page-length="{html.escape(self.options["data-page-length"])}"' if "data-page-length" in self.options else "" + csv_table_html = f'' head_rows = 0 head_rows += self.options.get('header-rows', 0) row_count = 0 diff --git a/docs/sphinx_setup/_static/css/custom.css b/docs/sphinx_setup/_static/css/custom.css index de8a05732a4d06..1679f7309da044 100644 --- a/docs/sphinx_setup/_static/css/custom.css +++ b/docs/sphinx_setup/_static/css/custom.css @@ -69,7 +69,7 @@ a#wap_dns { /* Sphinx-design tabs override */ .sd-tab-set>input:checked+label { color: var(--sd-color-black) !important; - background-color: #f8f8f8 !important; + background-color: white !important; border: solid 1px #bdbdbd; border-bottom: solid 0px; margin-bottom: -1px; @@ -96,7 +96,7 @@ a#wap_dns { cursor: pointer; font-size: var(--sd-fontsize-tabs-label); font-weight: 400 !important; - padding: 5px 16px 2px !important; + padding: 5px 16px 0px !important; transition: color 250ms; width: auto; z-index: 1; @@ -110,7 +110,6 @@ a#wap_dns { box-shadow: 0 0 0 0; border: solid 1px var(--sd-color-tabs-overline); border-color: #bdbdbd; - background-color: #f8f8f8; padding-right: 4px; padding-left: 4px; padding-bottom: 6px; diff --git a/docs/sphinx_setup/_static/css/openVinoDataTables.css b/docs/sphinx_setup/_static/css/openVinoDataTables.css index 526aabb6abe15d..bedc0f5206e260 100644 --- a/docs/sphinx_setup/_static/css/openVinoDataTables.css +++ b/docs/sphinx_setup/_static/css/openVinoDataTables.css @@ -6,8 +6,7 @@ div.dt-buttons>.dt-button, div.dt-buttons>div.dt-button-split .dt-button { } div.dt-container .dt-paging .dt-paging-button:hover { - color: white !important; - border: 1px solid #aaa; + border: 1px solid #aaa !important; background:none !important; background-color: var(--bttn-act-bg-hover) !important } @@ -190,10 +189,9 @@ div.dt-container .dt-paging .dt-paging-button { div.dt-container .dt-paging .dt-paging-button.current, div.dt-container .dt-paging .dt-paging-button.current:hover { background: none !important; - background-color: var(--bttn-act-bg-active) !important; + background-color: var(--bttn-sec-border-color) !important; border-color: var(--bttn-act-bg-active) !important; border-radius: 0px !important; - color: white !important; border: 1px !important } table.dataTable thead>tr>th.dt-orderable-asc span.dt-column-order:before, table.dataTable thead>tr>th.dt-orderable-asc span.dt-column-order:after, table.dataTable thead>tr>th.dt-orderable-desc span.dt-column-order:before, table.dataTable thead>tr>th.dt-orderable-desc span.dt-column-order:after, table.dataTable thead>tr>th.dt-ordering-asc span.dt-column-order:before, table.dataTable thead>tr>th.dt-ordering-asc span.dt-column-order:after, table.dataTable thead>tr>th.dt-ordering-desc span.dt-column-order:before, table.dataTable thead>tr>th.dt-ordering-desc span.dt-column-order:after, table.dataTable thead>tr>td.dt-orderable-asc span.dt-column-order:before, table.dataTable thead>tr>td.dt-orderable-asc span.dt-column-order:after, table.dataTable thead>tr>td.dt-orderable-desc span.dt-column-order:before, table.dataTable thead>tr>td.dt-orderable-desc span.dt-column-order:after, table.dataTable thead>tr>td.dt-ordering-asc span.dt-column-order:before, table.dataTable thead>tr>td.dt-ordering-asc span.dt-column-order:after, table.dataTable thead>tr>td.dt-ordering-desc span.dt-column-order:before, table.dataTable thead>tr>td.dt-ordering-desc span.dt-column-order:after { diff --git a/docs/sphinx_setup/_static/js/openVinoDataTables.js b/docs/sphinx_setup/_static/js/openVinoDataTables.js index bd56a71533786c..fb3a57d959020c 100644 --- a/docs/sphinx_setup/_static/js/openVinoDataTables.js +++ b/docs/sphinx_setup/_static/js/openVinoDataTables.js @@ -1,16 +1,15 @@ $(document).ready(function () { var columnDefs = []; - var tables = $('table.modeldata'); for (let table of tables) { - var hidden = table.getAttribute('data-columns-hidden'); + var hidden = table.getAttribute('data-column-hidden'); columnDefs = [{ "visible": false, "targets": JSON.parse(hidden) }] $(table).DataTable({ responsive: true, "autoWidth": false, language: { buttons: { - colvisRestore: "Restore default" + colvisRestore: "Restore default selection" } }, lengthMenu: [ diff --git a/docs/sphinx_setup/_templates/layout.html b/docs/sphinx_setup/_templates/layout.html index 0d2331b2c83fe3..a791091e1f13a4 100644 --- a/docs/sphinx_setup/_templates/layout.html +++ b/docs/sphinx_setup/_templates/layout.html @@ -9,6 +9,7 @@ + diff --git a/src/bindings/python/src/openvino/__init__.py b/src/bindings/python/src/openvino/__init__.py index 7643f742e0067d..69c678909b1c9e 100644 --- a/src/bindings/python/src/openvino/__init__.py +++ b/src/bindings/python/src/openvino/__init__.py @@ -7,7 +7,7 @@ # Required for Windows OS platforms # Note: always top-level try: - from openvino.package_utils import _add_openvino_libs_to_search_path + from openvino.utils import _add_openvino_libs_to_search_path _add_openvino_libs_to_search_path() except ImportError: pass @@ -17,47 +17,6 @@ # # This __init__.py forces checking of runtime modules to propagate errors. # # It is not compared with init files from openvino-dev package. # # - -# Openvino pybind bindings -from openvino._pyopenvino import AxisSet -from openvino._pyopenvino import AxisVector -from openvino._pyopenvino import ConstOutput -from openvino._pyopenvino import Coordinate -from openvino._pyopenvino import CoordinateDiff -from openvino._pyopenvino import DiscreteTypeInfo -from openvino._pyopenvino import Extension -from openvino._pyopenvino import ProfilingInfo -from openvino._pyopenvino import RTMap -from openvino._pyopenvino import Version -from openvino._pyopenvino import Symbol -from openvino._pyopenvino import Dimension -from openvino._pyopenvino import Input -from openvino._pyopenvino import Output -from openvino._pyopenvino import Node -from openvino._pyopenvino import Strides -from openvino._pyopenvino import PartialShape -from openvino._pyopenvino import Shape -from openvino._pyopenvino import Layout -from openvino._pyopenvino import Type -from openvino._pyopenvino import Tensor -from openvino._pyopenvino import OVAny -from openvino._pyopenvino import get_batch -from openvino._pyopenvino import set_batch -from openvino._pyopenvino import serialize -from openvino._pyopenvino import shutdown -from openvino._pyopenvino import save_model -from openvino._pyopenvino import layout_helpers -from openvino._pyopenvino import RemoteContext -from openvino._pyopenvino import RemoteTensor -from openvino._pyopenvino import Op - -# Import public classes from _ov_api -from openvino._ov_api import Model -from openvino._ov_api import Core -from openvino._ov_api import CompiledModel -from openvino._ov_api import InferRequest -from openvino._ov_api import AsyncInferQueue - # Import all public modules from openvino import runtime as runtime from openvino import frontend as frontend @@ -67,10 +26,36 @@ from openvino import utils as utils from openvino import properties as properties +# Import most important classes and functions from openvino.runtime +from openvino._ov_api import Model +from openvino._ov_api import Core +from openvino._ov_api import CompiledModel +from openvino._ov_api import InferRequest +from openvino._ov_api import AsyncInferQueue + +from openvino.runtime import Symbol +from openvino.runtime import Dimension +from openvino.runtime import Strides +from openvino.runtime import PartialShape +from openvino.runtime import Shape +from openvino.runtime import Layout +from openvino.runtime import Type +from openvino.runtime import Tensor +from openvino.runtime import OVAny + # Helper functions for openvino module -from openvino.utils.data_helpers import tensor_from_file +from openvino.runtime.utils.data_helpers import tensor_from_file from openvino._ov_api import compile_model +from openvino.runtime import get_batch +from openvino.runtime import set_batch +from openvino.runtime import serialize +from openvino.runtime import shutdown +from openvino.runtime import save_model +from openvino.runtime import layout_helpers +from openvino._pyopenvino import RemoteContext +from openvino._pyopenvino import RemoteTensor +from openvino._pyopenvino import Op # Import opsets from openvino import opset1 @@ -95,7 +80,7 @@ from openvino._pyopenvino import VASurfaceTensor # Set version for openvino package -from openvino._pyopenvino import get_version +from openvino.runtime import get_version __version__ = get_version() # Tools diff --git a/src/bindings/python/src/openvino/_ov_api.py b/src/bindings/python/src/openvino/_ov_api.py index da31fab4c95d8e..53d0fa5316498b 100644 --- a/src/bindings/python/src/openvino/_ov_api.py +++ b/src/bindings/python/src/openvino/_ov_api.py @@ -5,7 +5,9 @@ from types import TracebackType from typing import Any, Iterable, Union, Optional, Dict, Type from pathlib import Path +import warnings +import numpy as np from openvino._pyopenvino import Model as ModelBase from openvino._pyopenvino import Core as CoreBase @@ -14,7 +16,7 @@ from openvino._pyopenvino import Tensor from openvino._pyopenvino import Node -from openvino.utils.data_helpers import ( +from openvino.runtime.utils.data_helpers import ( OVDict, _InferRequestWrapper, _data_dispatch, diff --git a/src/bindings/python/src/openvino/frontend/frontend.py b/src/bindings/python/src/openvino/frontend/frontend.py index 6a16d5a573b7d7..4d549d24b4ef7c 100644 --- a/src/bindings/python/src/openvino/frontend/frontend.py +++ b/src/bindings/python/src/openvino/frontend/frontend.py @@ -7,7 +7,7 @@ from openvino._pyopenvino import FrontEnd as FrontEndBase from openvino._pyopenvino import FrontEndManager as FrontEndManagerBase from openvino._pyopenvino import InputModel -from openvino import Model +from openvino.runtime import Model class FrontEnd(FrontEndBase): diff --git a/src/bindings/python/src/openvino/frontend/jax/jaxpr_decoder.py b/src/bindings/python/src/openvino/frontend/jax/jaxpr_decoder.py index 9072598f824939..914f6b2e2ee548 100644 --- a/src/bindings/python/src/openvino/frontend/jax/jaxpr_decoder.py +++ b/src/bindings/python/src/openvino/frontend/jax/jaxpr_decoder.py @@ -6,7 +6,7 @@ import jax.core from openvino.frontend.jax.py_jax_frontend import _FrontEndJaxDecoder as Decoder -from openvino import PartialShape, Type as OVType, OVAny +from openvino.runtime import PartialShape, Type as OVType, OVAny from openvino.frontend.jax.utils import jax_array_to_ov_const, get_ov_type_for_value, \ ivalue_to_constant, param_to_constants diff --git a/src/bindings/python/src/openvino/frontend/jax/utils.py b/src/bindings/python/src/openvino/frontend/jax/utils.py index 659677b11d5af8..4535265d6de082 100644 --- a/src/bindings/python/src/openvino/frontend/jax/utils.py +++ b/src/bindings/python/src/openvino/frontend/jax/utils.py @@ -8,7 +8,7 @@ import jax.numpy as jnp import numpy as np from openvino.frontend.jax.passes import filter_element, filter_ivalue, filter_param -from openvino import op, Type as OVType, Shape, OVAny +from openvino.runtime import op, Type as OVType, Shape, OVAny numpy_to_ov_type_map = { np.float32: OVType.f32, diff --git a/src/bindings/python/src/openvino/frontend/pytorch/fx_decoder.py b/src/bindings/python/src/openvino/frontend/pytorch/fx_decoder.py index 81a2764ee1188d..c448571f1ac17a 100644 --- a/src/bindings/python/src/openvino/frontend/pytorch/fx_decoder.py +++ b/src/bindings/python/src/openvino/frontend/pytorch/fx_decoder.py @@ -10,7 +10,7 @@ from openvino.frontend.pytorch.py_pytorch_frontend import _FrontEndPytorchDecoder as Decoder from openvino.frontend.pytorch.py_pytorch_frontend import _Type as DecoderType -from openvino import PartialShape, Type as OVType, OVAny, Shape +from openvino.runtime import PartialShape, Type as OVType, OVAny, Shape from openvino.frontend.pytorch.utils import make_constant, fetch_attr, pt_to_ov_type_map, torch_tensor_to_ov_const logger = logging.getLogger(__name__) diff --git a/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/backend.py b/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/backend.py index a9a65781dcb254..9f2ef019769875 100644 --- a/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/backend.py +++ b/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/backend.py @@ -18,7 +18,7 @@ from torch._decomp import decomposition_table, get_decompositions from openvino.frontend import FrontEndManager -from openvino import Core, Type, PartialShape +from openvino.runtime import Core, Type, PartialShape from openvino.frontend.pytorch.ts_decoder import TorchScriptPythonDecoder from openvino.frontend.pytorch.torchdynamo import decompositions from openvino.frontend.pytorch.torchdynamo.decompositions import get_aot_decomposition_list, get_inf_decomposition_list @@ -27,7 +27,7 @@ from openvino.frontend.pytorch.torchdynamo.compile import cached_model_name, openvino_compile_cached_model from openvino.frontend.pytorch.torchdynamo.backend_utils import _get_cache_dir, _get_device, _get_model_caching, _get_decompositions, _get_aot_autograd -from openvino import Core, Type, PartialShape +from openvino.runtime import Core, Type, PartialShape logger = logging.getLogger(__name__) logger.setLevel(logging.WARNING) diff --git a/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/backend_utils.py b/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/backend_utils.py index c9a772b3feac42..47b3b82806b18b 100644 --- a/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/backend_utils.py +++ b/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/backend_utils.py @@ -5,7 +5,7 @@ # mypy: ignore-errors from typing import Optional, Any -from openvino import Core +from openvino.runtime import Core def _get_device(options) -> Optional[Any]: diff --git a/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/compile.py b/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/compile.py index ca8d5478e76c15..fa446893a05d07 100644 --- a/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/compile.py +++ b/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/compile.py @@ -14,7 +14,7 @@ from openvino.frontend import FrontEndManager from openvino.frontend.pytorch.fx_decoder import TorchFXPythonDecoder -from openvino import Core, Type, PartialShape, serialize +from openvino.runtime import Core, Type, PartialShape, serialize from openvino.frontend.pytorch.torchdynamo.backend_utils import _get_cache_dir, _get_device, _get_config, _is_cache_dir_in_config from typing import Callable, Optional diff --git a/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/execute.py b/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/execute.py index 7527ad7acb37a4..4f41f7b5a6a9de 100644 --- a/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/execute.py +++ b/src/bindings/python/src/openvino/frontend/pytorch/torchdynamo/execute.py @@ -20,7 +20,7 @@ from openvino.frontend.pytorch.fx_decoder import TorchFXPythonDecoder from openvino.frontend.pytorch.torchdynamo.partition import Partitioner from openvino.frontend.pytorch.torchdynamo.compile import openvino_compile -from openvino import Core, Type, PartialShape +from openvino.runtime import Core, Type, PartialShape from openvino.frontend.pytorch.torchdynamo.backend_utils import _get_cache_dir, _get_device, _get_aot_autograd from typing import Callable, Optional, Any diff --git a/src/bindings/python/src/openvino/frontend/pytorch/ts_decoder.py b/src/bindings/python/src/openvino/frontend/pytorch/ts_decoder.py index 7bb8073167a654..6d8fdb1658793e 100644 --- a/src/bindings/python/src/openvino/frontend/pytorch/ts_decoder.py +++ b/src/bindings/python/src/openvino/frontend/pytorch/ts_decoder.py @@ -6,7 +6,7 @@ from openvino.frontend.pytorch.py_pytorch_frontend import _FrontEndPytorchDecoder as Decoder from openvino.frontend.pytorch.py_pytorch_frontend import _Type as DecoderType -from openvino import op, PartialShape, Type as OVType, OVAny +from openvino.runtime import op, PartialShape, Type as OVType, OVAny from openvino.frontend.pytorch.utils import ( ivalue_to_constant, get_value_from_getattr, @@ -15,7 +15,7 @@ convert_quantized_tensor, graph_has_ops, ) -from openvino import opset11 as ops +from openvino.runtime import opset11 as ops from openvino.frontend.pytorch import quantized, patch_model from openvino.frontend.pytorch.module_extension import ModuleExtension diff --git a/src/bindings/python/src/openvino/frontend/pytorch/utils.py b/src/bindings/python/src/openvino/frontend/pytorch/utils.py index 9ba36707037c9e..826d766505fa79 100644 --- a/src/bindings/python/src/openvino/frontend/pytorch/utils.py +++ b/src/bindings/python/src/openvino/frontend/pytorch/utils.py @@ -7,8 +7,8 @@ import torch import numpy as np -from openvino import op, Type as OVType, Shape, Tensor -from openvino import opset11 as ops +from openvino.runtime import op, Type as OVType, Shape, Tensor +from openvino.runtime import opset11 as ops def make_constant(*args, **kwargs): diff --git a/src/bindings/python/src/openvino/frontend/tensorflow/node_decoder.py b/src/bindings/python/src/openvino/frontend/tensorflow/node_decoder.py index d15262cbc30366..fcedd7a74c2b51 100644 --- a/src/bindings/python/src/openvino/frontend/tensorflow/node_decoder.py +++ b/src/bindings/python/src/openvino/frontend/tensorflow/node_decoder.py @@ -7,7 +7,7 @@ import numpy as np import tensorflow as tf from openvino.frontend.tensorflow.py_tensorflow_frontend import _FrontEndDecoderBase as DecoderBase -from openvino import PartialShape, Type, OVAny, Tensor +from openvino.runtime import PartialShape, Type, OVAny, Tensor def tf_type_to_ov_type(tf_type_int): diff --git a/src/bindings/python/src/openvino/frontend/tensorflow/utils.py b/src/bindings/python/src/openvino/frontend/tensorflow/utils.py index 7de5dc950be53e..74c0dfff92297e 100644 --- a/src/bindings/python/src/openvino/frontend/tensorflow/utils.py +++ b/src/bindings/python/src/openvino/frontend/tensorflow/utils.py @@ -8,7 +8,7 @@ import logging as log import numpy as np import sys -from openvino import PartialShape, Dimension, Type +from openvino.runtime import PartialShape, Dimension, Type from packaging.version import parse, Version from typing import List, Dict, Union diff --git a/src/bindings/python/src/openvino/helpers/packing.py b/src/bindings/python/src/openvino/helpers/packing.py index d0956e09fc6261..796af87402f3a6 100644 --- a/src/bindings/python/src/openvino/helpers/packing.py +++ b/src/bindings/python/src/openvino/helpers/packing.py @@ -5,7 +5,7 @@ import numpy as np from typing import Union -from openvino import Type, Shape +from openvino.runtime import Type, Shape def pack_data(array: np.ndarray, type: Type) -> np.ndarray: diff --git a/src/bindings/python/src/openvino/opset1/ops.py b/src/bindings/python/src/openvino/opset1/ops.py index e264aea304fb1f..edca6c62a0b246 100644 --- a/src/bindings/python/src/openvino/opset1/ops.py +++ b/src/bindings/python/src/openvino/opset1/ops.py @@ -8,17 +8,17 @@ import numpy as np from functools import partial -from openvino import Node, PartialShape, Type +from openvino.runtime import Node, PartialShape, Type from openvino.op import Constant, Parameter, tensor_iterator -from openvino.utils.node_factory import _get_node_factory -from openvino.utils.decorators import binary_op, nameable_op, unary_op -from openvino.utils.input_validation import ( +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import binary_op, nameable_op, unary_op +from openvino.runtime.utils.input_validation import ( check_valid_attributes, is_non_negative_value, is_positive_value, ) -from openvino.utils.node_factory import NodeFactory -from openvino.utils.types import ( +from openvino.runtime.utils.node_factory import NodeFactory +from openvino.runtime.utils.types import ( NodeInput, NumericData, NumericType, diff --git a/src/bindings/python/src/openvino/opset10/ops.py b/src/bindings/python/src/openvino/opset10/ops.py index d0bc3cbf1cba4a..c7b75777484a59 100644 --- a/src/bindings/python/src/openvino/opset10/ops.py +++ b/src/bindings/python/src/openvino/opset10/ops.py @@ -6,10 +6,10 @@ from functools import partial from typing import List, Optional -from openvino import Node -from openvino.utils.node_factory import _get_node_factory -from openvino.utils.decorators import nameable_op -from openvino.utils.types import ( +from openvino.runtime import Node +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import nameable_op +from openvino.runtime.utils.types import ( NodeInput, as_nodes, as_node, diff --git a/src/bindings/python/src/openvino/opset11/ops.py b/src/bindings/python/src/openvino/opset11/ops.py index 95767b4800db1c..575c99501d2d6c 100644 --- a/src/bindings/python/src/openvino/opset11/ops.py +++ b/src/bindings/python/src/openvino/opset11/ops.py @@ -6,10 +6,10 @@ from functools import partial from typing import List, Optional -from openvino import Node -from openvino.utils.node_factory import _get_node_factory -from openvino.utils.decorators import nameable_op -from openvino.utils.types import ( +from openvino.runtime import Node +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import nameable_op +from openvino.runtime.utils.types import ( NodeInput, as_nodes, ) diff --git a/src/bindings/python/src/openvino/opset12/ops.py b/src/bindings/python/src/openvino/opset12/ops.py index 4b354b1fcff973..928bf4f71a9773 100644 --- a/src/bindings/python/src/openvino/opset12/ops.py +++ b/src/bindings/python/src/openvino/opset12/ops.py @@ -6,10 +6,10 @@ from functools import partial from typing import Optional -from openvino import Node -from openvino.utils.node_factory import _get_node_factory -from openvino.utils.decorators import nameable_op -from openvino.utils.types import ( +from openvino.runtime import Node +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import nameable_op +from openvino.runtime.utils.types import ( NodeInput, as_nodes, as_node, diff --git a/src/bindings/python/src/openvino/opset13/ops.py b/src/bindings/python/src/openvino/opset13/ops.py index 5c6863740120f8..12f0d06b1a28e6 100644 --- a/src/bindings/python/src/openvino/opset13/ops.py +++ b/src/bindings/python/src/openvino/opset13/ops.py @@ -11,12 +11,12 @@ log = logging.getLogger(__name__) -from openvino import Node, Shape, Type, Output, Tensor +from openvino.runtime import Node, Shape, Type, Output, Tensor from openvino.op import Constant, Result from openvino.opset1 import convert_like -from openvino.utils.node_factory import _get_node_factory -from openvino.utils.decorators import binary_op, nameable_op, unary_op, overloading -from openvino.utils.types import ( +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import binary_op, nameable_op, unary_op, overloading +from openvino.runtime.utils.types import ( NumericData, NodeInput, NumericType, diff --git a/src/bindings/python/src/openvino/opset14/ops.py b/src/bindings/python/src/openvino/opset14/ops.py index 59e1bfd3e89c6f..fa872d24eb7f1a 100644 --- a/src/bindings/python/src/openvino/opset14/ops.py +++ b/src/bindings/python/src/openvino/opset14/ops.py @@ -7,11 +7,11 @@ from typing import Union, Optional, List -from openvino import Node, Type -from openvino.utils.node_factory import _get_node_factory -from openvino.utils.types import TensorShape -from openvino.utils.decorators import nameable_op -from openvino.utils.types import NodeInput, as_node, as_nodes +from openvino.runtime import Node, Type +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.types import TensorShape +from openvino.runtime.utils.decorators import nameable_op +from openvino.runtime.utils.types import NodeInput, as_node, as_nodes _get_node_factory_opset14 = partial(_get_node_factory, "opset14") diff --git a/src/bindings/python/src/openvino/opset15/ops.py b/src/bindings/python/src/openvino/opset15/ops.py index 97d4419fc4834b..8e6b8bd46d5f7c 100644 --- a/src/bindings/python/src/openvino/opset15/ops.py +++ b/src/bindings/python/src/openvino/opset15/ops.py @@ -7,12 +7,12 @@ from typing import List, Literal, Optional import numpy as np -from openvino import Node, Type +from openvino.runtime import Node, Type from openvino.opset1 import convert_like from openvino.opset14 import constant -from openvino.utils.node_factory import _get_node_factory -from openvino.utils.decorators import binary_op, nameable_op -from openvino.utils.types import NodeInput, as_nodes +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import binary_op, nameable_op +from openvino.runtime.utils.types import NodeInput, as_nodes _get_node_factory_opset15 = partial(_get_node_factory, "opset15") diff --git a/src/bindings/python/src/openvino/opset16/ops.py b/src/bindings/python/src/openvino/opset16/ops.py index e5ebdc7a2a11d6..60656f6d993b6a 100644 --- a/src/bindings/python/src/openvino/opset16/ops.py +++ b/src/bindings/python/src/openvino/opset16/ops.py @@ -6,10 +6,10 @@ from functools import partial from typing import Optional -from openvino import Node -from openvino.utils.decorators import nameable_op -from openvino.utils.node_factory import _get_node_factory -from openvino.utils.types import NodeInput, as_nodes +from openvino.runtime import Node +from openvino.runtime.utils.decorators import nameable_op +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.types import NodeInput, as_nodes _get_node_factory_opset16 = partial(_get_node_factory, "opset16") diff --git a/src/bindings/python/src/openvino/opset2/ops.py b/src/bindings/python/src/openvino/opset2/ops.py index f76f608fe9a5c7..45b33f5bc0288b 100644 --- a/src/bindings/python/src/openvino/opset2/ops.py +++ b/src/bindings/python/src/openvino/opset2/ops.py @@ -9,17 +9,18 @@ from functools import partial import warnings -from openvino import Node, Shape +from openvino.runtime import Node, Shape from openvino.op import Constant, Parameter -from openvino.utils.decorators import binary_op, nameable_op, unary_op -from openvino.utils.input_validation import ( +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import binary_op, nameable_op, unary_op +from openvino.runtime.utils.input_validation import ( assert_list_of_ints, check_valid_attributes, is_non_negative_value, is_positive_value, ) -from openvino.utils.node_factory import NodeFactory, _get_node_factory -from openvino.utils.types import ( +from openvino.runtime.utils.node_factory import NodeFactory +from openvino.runtime.utils.types import ( NodeInput, NumericData, NumericType, diff --git a/src/bindings/python/src/openvino/opset3/ops.py b/src/bindings/python/src/openvino/opset3/ops.py index 1c2c7e309fe919..989f5819acb685 100644 --- a/src/bindings/python/src/openvino/opset3/ops.py +++ b/src/bindings/python/src/openvino/opset3/ops.py @@ -8,17 +8,18 @@ import numpy as np from functools import partial -from openvino import Node, Shape +from openvino.runtime import Node, Shape from openvino.op import Constant, Parameter -from openvino.utils.decorators import binary_op, nameable_op, unary_op -from openvino.utils.input_validation import ( +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import binary_op, nameable_op, unary_op +from openvino.runtime.utils.input_validation import ( assert_list_of_ints, check_valid_attributes, is_non_negative_value, is_positive_value, ) -from openvino.utils.node_factory import NodeFactory, _get_node_factory -from openvino.utils.types import ( +from openvino.runtime.utils.node_factory import NodeFactory +from openvino.runtime.utils.types import ( NodeInput, NumericData, NumericType, diff --git a/src/bindings/python/src/openvino/opset4/ops.py b/src/bindings/python/src/openvino/opset4/ops.py index e6f3a3a1550937..4f6ba016852b02 100644 --- a/src/bindings/python/src/openvino/opset4/ops.py +++ b/src/bindings/python/src/openvino/opset4/ops.py @@ -8,17 +8,18 @@ import numpy as np from functools import partial -from openvino import Node, Shape +from openvino.runtime import Node, Shape from openvino.op import Constant, Parameter -from openvino.utils.decorators import binary_op, nameable_op, unary_op -from openvino.utils.input_validation import ( +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import binary_op, nameable_op, unary_op +from openvino.runtime.utils.input_validation import ( assert_list_of_ints, check_valid_attributes, is_non_negative_value, is_positive_value, ) -from openvino.utils.node_factory import NodeFactory, _get_node_factory -from openvino.utils.types import ( +from openvino.runtime.utils.node_factory import NodeFactory +from openvino.runtime.utils.types import ( NodeInput, NumericData, NumericType, diff --git a/src/bindings/python/src/openvino/opset5/ops.py b/src/bindings/python/src/openvino/opset5/ops.py index 9217830752b1d8..20057b78c7c31d 100644 --- a/src/bindings/python/src/openvino/opset5/ops.py +++ b/src/bindings/python/src/openvino/opset5/ops.py @@ -8,17 +8,18 @@ import numpy as np from functools import partial -from openvino import Node, Shape +from openvino.runtime import Node, Shape from openvino.op import Constant, Parameter, loop -from openvino.utils.decorators import binary_op, nameable_op, unary_op -from openvino.utils.input_validation import ( +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import binary_op, nameable_op, unary_op +from openvino.runtime.utils.input_validation import ( assert_list_of_ints, check_valid_attributes, is_non_negative_value, is_positive_value, ) -from openvino.utils.node_factory import NodeFactory, _get_node_factory -from openvino.utils.types import ( +from openvino.runtime.utils.node_factory import NodeFactory +from openvino.runtime.utils.types import ( NodeInput, NumericData, NumericType, diff --git a/src/bindings/python/src/openvino/opset6/ops.py b/src/bindings/python/src/openvino/opset6/ops.py index 340d0405b4ba23..8020715f20dea3 100644 --- a/src/bindings/python/src/openvino/opset6/ops.py +++ b/src/bindings/python/src/openvino/opset6/ops.py @@ -9,13 +9,13 @@ from functools import partial, singledispatch -from openvino import Node, Type, PartialShape, Output, Shape +from openvino.runtime import Node, Type, PartialShape, Output, Shape from openvino.op import assign, Constant, Parameter from openvino.op import read_value as _read_value from openvino.op.util import VariableInfo, Variable -from openvino.utils.node_factory import _get_node_factory -from openvino.utils.decorators import nameable_op, overloading -from openvino.utils.types import ( +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import nameable_op, overloading +from openvino.runtime.utils.types import ( NodeInput, NumericType, TensorShape, diff --git a/src/bindings/python/src/openvino/opset7/ops.py b/src/bindings/python/src/openvino/opset7/ops.py index e33d266debedf1..59e09b64888eb1 100644 --- a/src/bindings/python/src/openvino/opset7/ops.py +++ b/src/bindings/python/src/openvino/opset7/ops.py @@ -7,17 +7,18 @@ from typing import Callable, Iterable, List, Optional, Set, Union import numpy as np -from openvino import Node, Shape +from openvino.runtime import Node, Shape from openvino.op import Constant, Parameter -from openvino.utils.decorators import binary_op, nameable_op, unary_op -from openvino.utils.input_validation import ( +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import binary_op, nameable_op, unary_op +from openvino.runtime.utils.input_validation import ( assert_list_of_ints, check_valid_attributes, is_non_negative_value, is_positive_value, ) -from openvino.utils.node_factory import NodeFactory, _get_node_factory -from openvino.utils.types import ( +from openvino.runtime.utils.node_factory import NodeFactory +from openvino.runtime.utils.types import ( NodeInput, NumericData, NumericType, diff --git a/src/bindings/python/src/openvino/opset8/ops.py b/src/bindings/python/src/openvino/opset8/ops.py index a9a868e7b541d8..6995d55a28a776 100644 --- a/src/bindings/python/src/openvino/opset8/ops.py +++ b/src/bindings/python/src/openvino/opset8/ops.py @@ -9,15 +9,15 @@ import numpy as np from openvino.exceptions import UserInputError from openvino.op import Constant, Parameter, if_op -from openvino import Node -from openvino.utils.node_factory import _get_node_factory -from openvino.utils.decorators import nameable_op -from openvino.utils.input_validation import ( +from openvino.runtime import Node +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import nameable_op +from openvino.runtime.utils.input_validation import ( check_valid_attributes, is_non_negative_value, is_positive_value, ) -from openvino.utils.types import ( +from openvino.runtime.utils.types import ( NodeInput, TensorShape, as_node, diff --git a/src/bindings/python/src/openvino/opset9/ops.py b/src/bindings/python/src/openvino/opset9/ops.py index e2264845e058dc..a6d45cfd0be2cc 100644 --- a/src/bindings/python/src/openvino/opset9/ops.py +++ b/src/bindings/python/src/openvino/opset9/ops.py @@ -7,10 +7,10 @@ from typing import Optional import numpy as np -from openvino import Node -from openvino.utils.node_factory import _get_node_factory -from openvino.utils.decorators import nameable_op -from openvino.utils.types import ( +from openvino.runtime import Node +from openvino.runtime.opset_utils import _get_node_factory +from openvino.runtime.utils.decorators import nameable_op +from openvino.runtime.utils.types import ( NodeInput, as_nodes, as_node, diff --git a/src/bindings/python/src/openvino/preprocess/torchvision/preprocess_converter.py b/src/bindings/python/src/openvino/preprocess/torchvision/preprocess_converter.py index 717e945217468c..c14635cc118208 100644 --- a/src/bindings/python/src/openvino/preprocess/torchvision/preprocess_converter.py +++ b/src/bindings/python/src/openvino/preprocess/torchvision/preprocess_converter.py @@ -5,7 +5,7 @@ from typing import Callable, Any, Union import logging -import openvino as ov +import openvino.runtime as ov class PreprocessConverter(): diff --git a/src/bindings/python/src/openvino/preprocess/torchvision/torchvision_preprocessing.py b/src/bindings/python/src/openvino/preprocess/torchvision/torchvision_preprocessing.py index 5dad42b47da44a..f8b51afd546f57 100644 --- a/src/bindings/python/src/openvino/preprocess/torchvision/torchvision_preprocessing.py +++ b/src/bindings/python/src/openvino/preprocess/torchvision/torchvision_preprocessing.py @@ -20,10 +20,10 @@ import torchvision.transforms as transforms from torchvision.transforms import InterpolationMode -import openvino as ov -import openvino.opset11 as ops -from openvino import Layout, Type -from openvino.utils.decorators import custom_preprocess_function +import openvino.runtime as ov +import openvino.runtime.opset11 as ops +from openvino.runtime import Layout, Type +from openvino.runtime.utils.decorators import custom_preprocess_function from openvino.preprocess import PrePostProcessor, ResizeAlgorithm, ColorFormat diff --git a/src/bindings/python/src/openvino/runtime/opset_utils.py b/src/bindings/python/src/openvino/runtime/opset_utils.py new file mode 100644 index 00000000000000..475750e71f87c5 --- /dev/null +++ b/src/bindings/python/src/openvino/runtime/opset_utils.py @@ -0,0 +1,22 @@ +# -*- coding: utf-8 -*- +# Copyright (C) 2018-2024 Intel Corporation +# SPDX-License-Identifier: Apache-2.0 + +from typing import Optional +import numpy as np + +from openvino.runtime import Node +from openvino.runtime.utils.decorators import nameable_op +from openvino.runtime.utils.node_factory import NodeFactory +from openvino.runtime.utils.types import ( + as_node, + NodeInput, +) + + +def _get_node_factory(opset_version: Optional[str] = None) -> NodeFactory: + """Return NodeFactory configured to create operators from specified opset version.""" + if opset_version: + return NodeFactory(opset_version) + else: + return NodeFactory() diff --git a/src/bindings/python/src/openvino/runtime/opset_utils/__init__.py b/src/bindings/python/src/openvino/runtime/opset_utils/__init__.py deleted file mode 100644 index 6fb3e5f6f0c950..00000000000000 --- a/src/bindings/python/src/openvino/runtime/opset_utils/__init__.py +++ /dev/null @@ -1,6 +0,0 @@ -# -*- coding: utf-8 -*- -# Copyright (C) 2018-2024 Intel Corporation -# SPDX-License-Identifier: Apache-2.0 - - -from openvino.utils.node_factory import _get_node_factory diff --git a/src/bindings/python/src/openvino/runtime/utils/__init__.py b/src/bindings/python/src/openvino/runtime/utils/__init__.py index 8447e93a907277..73399ccbed2598 100644 --- a/src/bindings/python/src/openvino/runtime/utils/__init__.py +++ b/src/bindings/python/src/openvino/runtime/utils/__init__.py @@ -4,4 +4,4 @@ """Generic utilities. Factor related functions out to separate files.""" -from openvino.utils import numpy_to_c, replace_node, replace_output_update_name +from openvino._pyopenvino.util import numpy_to_c, replace_node, replace_output_update_name diff --git a/src/bindings/python/src/openvino/utils/broadcasting.py b/src/bindings/python/src/openvino/runtime/utils/broadcasting.py similarity index 87% rename from src/bindings/python/src/openvino/utils/broadcasting.py rename to src/bindings/python/src/openvino/runtime/utils/broadcasting.py index 01549625e2c628..9fd13da7728e29 100644 --- a/src/bindings/python/src/openvino/utils/broadcasting.py +++ b/src/bindings/python/src/openvino/runtime/utils/broadcasting.py @@ -3,11 +3,14 @@ # SPDX-License-Identifier: Apache-2.0 import logging -from typing import Optional +from typing import List, Optional -from openvino import AxisSet -from openvino.utils.types import ( +from openvino.runtime import AxisSet, Node +from openvino.runtime.utils.types import ( + NodeInput, TensorShape, + get_dtype, + make_constant_node, ) log = logging.getLogger(__name__) diff --git a/src/bindings/python/src/openvino/runtime/utils/broadcasting/__init__.py b/src/bindings/python/src/openvino/runtime/utils/broadcasting/__init__.py deleted file mode 100644 index 3219f239f0ab44..00000000000000 --- a/src/bindings/python/src/openvino/runtime/utils/broadcasting/__init__.py +++ /dev/null @@ -1,5 +0,0 @@ -# -*- coding: utf-8 -*- -# Copyright (C) 2018-2024 Intel Corporation -# SPDX-License-Identifier: Apache-2.0 - -from openvino.utils.broadcasting import get_broadcast_axes diff --git a/src/bindings/python/src/openvino/runtime/utils/data_helpers/__init__.py b/src/bindings/python/src/openvino/runtime/utils/data_helpers/__init__.py index 282547dd9df79a..a46105efaaeadb 100644 --- a/src/bindings/python/src/openvino/runtime/utils/data_helpers/__init__.py +++ b/src/bindings/python/src/openvino/runtime/utils/data_helpers/__init__.py @@ -2,7 +2,7 @@ # Copyright (C) 2018-2024 Intel Corporation # SPDX-License-Identifier: Apache-2.0 -from openvino.utils.data_helpers.data_dispatcher import _data_dispatch -from openvino.utils.data_helpers.wrappers import tensor_from_file -from openvino.utils.data_helpers.wrappers import _InferRequestWrapper -from openvino.utils.data_helpers.wrappers import OVDict +from openvino.runtime.utils.data_helpers.data_dispatcher import _data_dispatch +from openvino.runtime.utils.data_helpers.wrappers import tensor_from_file +from openvino.runtime.utils.data_helpers.wrappers import _InferRequestWrapper +from openvino.runtime.utils.data_helpers.wrappers import OVDict diff --git a/src/bindings/python/src/openvino/utils/data_helpers/data_dispatcher.py b/src/bindings/python/src/openvino/runtime/utils/data_helpers/data_dispatcher.py similarity index 99% rename from src/bindings/python/src/openvino/utils/data_helpers/data_dispatcher.py rename to src/bindings/python/src/openvino/runtime/utils/data_helpers/data_dispatcher.py index d4db7cb07b629c..bce10c9c3774ef 100644 --- a/src/bindings/python/src/openvino/utils/data_helpers/data_dispatcher.py +++ b/src/bindings/python/src/openvino/runtime/utils/data_helpers/data_dispatcher.py @@ -8,7 +8,7 @@ import numpy as np from openvino._pyopenvino import ConstOutput, Tensor, Type, RemoteTensor -from openvino.utils.data_helpers.wrappers import _InferRequestWrapper, OVDict +from openvino.runtime.utils.data_helpers.wrappers import _InferRequestWrapper, OVDict ContainerTypes = Union[dict, list, tuple, OVDict] ScalarTypes = Union[np.number, int, float] diff --git a/src/bindings/python/src/openvino/runtime/utils/data_helpers/data_dispatcher/__init__.py b/src/bindings/python/src/openvino/runtime/utils/data_helpers/data_dispatcher/__init__.py deleted file mode 100644 index e0a2d022660dd3..00000000000000 --- a/src/bindings/python/src/openvino/runtime/utils/data_helpers/data_dispatcher/__init__.py +++ /dev/null @@ -1,20 +0,0 @@ -# -*- coding: utf-8 -*- -# Copyright (C) 2018-2024 Intel Corporation -# SPDX-License-Identifier: Apache-2.0 - - -from openvino.utils.data_helpers.data_dispatcher import ContainerTypes -from openvino.utils.data_helpers.data_dispatcher import ScalarTypes -from openvino.utils.data_helpers.data_dispatcher import ValidKeys - -from openvino.utils.data_helpers.data_dispatcher import is_list_simple_type -from openvino.utils.data_helpers.data_dispatcher import get_request_tensor -from openvino.utils.data_helpers.data_dispatcher import value_to_tensor -from openvino.utils.data_helpers.data_dispatcher import to_c_style -from openvino.utils.data_helpers.data_dispatcher import normalize_arrays -from openvino.utils.data_helpers.data_dispatcher import create_shared -from openvino.utils.data_helpers.data_dispatcher import set_request_tensor -from openvino.utils.data_helpers.data_dispatcher import update_tensor -from openvino.utils.data_helpers.data_dispatcher import update_inputs -from openvino.utils.data_helpers.data_dispatcher import create_copied -from openvino.utils.data_helpers.data_dispatcher import _data_dispatch diff --git a/src/bindings/python/src/openvino/utils/data_helpers/wrappers.py b/src/bindings/python/src/openvino/runtime/utils/data_helpers/wrappers.py similarity index 100% rename from src/bindings/python/src/openvino/utils/data_helpers/wrappers.py rename to src/bindings/python/src/openvino/runtime/utils/data_helpers/wrappers.py diff --git a/src/bindings/python/src/openvino/runtime/utils/data_helpers/wrappers/__init__.py b/src/bindings/python/src/openvino/runtime/utils/data_helpers/wrappers/__init__.py deleted file mode 100644 index 22214fd24682da..00000000000000 --- a/src/bindings/python/src/openvino/runtime/utils/data_helpers/wrappers/__init__.py +++ /dev/null @@ -1,8 +0,0 @@ -# -*- coding: utf-8 -*- -# Copyright (C) 2018-2024 Intel Corporation -# SPDX-License-Identifier: Apache-2.0 - - -from openvino.utils.data_helpers.wrappers import tensor_from_file -from openvino.utils.data_helpers.wrappers import _InferRequestWrapper -from openvino.utils.data_helpers.wrappers import OVDict diff --git a/src/bindings/python/src/openvino/utils/decorators.py b/src/bindings/python/src/openvino/runtime/utils/decorators.py similarity index 98% rename from src/bindings/python/src/openvino/utils/decorators.py rename to src/bindings/python/src/openvino/runtime/utils/decorators.py index 9418c359d129e8..98da1ba4389ef7 100644 --- a/src/bindings/python/src/openvino/utils/decorators.py +++ b/src/bindings/python/src/openvino/runtime/utils/decorators.py @@ -6,8 +6,8 @@ from inspect import signature from typing import Any, Callable, Dict, Optional, Union, get_origin, get_args -from openvino import Node, Output -from openvino.utils.types import NodeInput, as_node, as_nodes +from openvino.runtime import Node, Output +from openvino.runtime.utils.types import NodeInput, as_node, as_nodes def _get_name(**kwargs: Any) -> Node: diff --git a/src/bindings/python/src/openvino/runtime/utils/decorators/__init__.py b/src/bindings/python/src/openvino/runtime/utils/decorators/__init__.py deleted file mode 100644 index bb0bac112d2c5f..00000000000000 --- a/src/bindings/python/src/openvino/runtime/utils/decorators/__init__.py +++ /dev/null @@ -1,13 +0,0 @@ -# -*- coding: utf-8 -*- -# Copyright (C) 2018-2024 Intel Corporation -# SPDX-License-Identifier: Apache-2.0 - -from openvino.utils.decorators import _get_name -from openvino.utils.decorators import _set_node_friendly_name -from openvino.utils.decorators import nameable_op -from openvino.utils.decorators import unary_op -from openvino.utils.decorators import binary_op -from openvino.utils.decorators import custom_preprocess_function -from openvino.utils.decorators import MultiMethod -from openvino.utils.decorators import registry -from openvino.utils.decorators import overloading diff --git a/src/bindings/python/src/openvino/utils/input_validation.py b/src/bindings/python/src/openvino/runtime/utils/input_validation.py similarity index 98% rename from src/bindings/python/src/openvino/utils/input_validation.py rename to src/bindings/python/src/openvino/runtime/utils/input_validation.py index 1de08452e1da9f..e79a16c48581b1 100644 --- a/src/bindings/python/src/openvino/utils/input_validation.py +++ b/src/bindings/python/src/openvino/runtime/utils/input_validation.py @@ -9,7 +9,7 @@ import numpy as np -from openvino.exceptions import UserInputError +from openvino.runtime.exceptions import UserInputError log = logging.getLogger(__name__) diff --git a/src/bindings/python/src/openvino/runtime/utils/input_validation/__init__.py b/src/bindings/python/src/openvino/runtime/utils/input_validation/__init__.py deleted file mode 100644 index 0b49e9ea33c40d..00000000000000 --- a/src/bindings/python/src/openvino/runtime/utils/input_validation/__init__.py +++ /dev/null @@ -1,10 +0,0 @@ -# -*- coding: utf-8 -*- -# Copyright (C) 2018-2024 Intel Corporation -# SPDX-License-Identifier: Apache-2.0 - -from openvino.utils.input_validation import assert_list_of_ints -from openvino.utils.input_validation import _check_value -from openvino.utils.input_validation import check_valid_attribute -from openvino.utils.input_validation import check_valid_attributes -from openvino.utils.input_validation import is_positive_value -from openvino.utils.input_validation import is_non_negative_value diff --git a/src/bindings/python/src/openvino/utils/node_factory.py b/src/bindings/python/src/openvino/runtime/utils/node_factory.py similarity index 92% rename from src/bindings/python/src/openvino/utils/node_factory.py rename to src/bindings/python/src/openvino/runtime/utils/node_factory.py index e999ae6988814a..25daf739223dba 100644 --- a/src/bindings/python/src/openvino/utils/node_factory.py +++ b/src/bindings/python/src/openvino/runtime/utils/node_factory.py @@ -2,16 +2,17 @@ # Copyright (C) 2018-2024 Intel Corporation # SPDX-License-Identifier: Apache-2.0 +import logging as log -from functools import singledispatchmethod +from functools import partial, singledispatchmethod from typing import Any, Dict, List, Optional, Union from pathlib import Path from openvino._pyopenvino import NodeFactory as _NodeFactory -from openvino import Node, Output, Extension +from openvino.runtime import Node, Output, Extension -from openvino.exceptions import UserInputError +from openvino.runtime.exceptions import UserInputError DEFAULT_OPSET = "opset13" @@ -124,11 +125,3 @@ def _arguments_as_outputs(arguments: List[Union[Node, Output]]) -> List[Output]: else: outputs.extend(argument.outputs()) return outputs - - -def _get_node_factory(opset_version: Optional[str] = None) -> NodeFactory: - """Return NodeFactory configured to create operators from specified opset version.""" - if opset_version: - return NodeFactory(opset_version) - else: - return NodeFactory() diff --git a/src/bindings/python/src/openvino/runtime/utils/node_factory/__init__.py b/src/bindings/python/src/openvino/runtime/utils/node_factory/__init__.py deleted file mode 100644 index 945ea8deb7863c..00000000000000 --- a/src/bindings/python/src/openvino/runtime/utils/node_factory/__init__.py +++ /dev/null @@ -1,5 +0,0 @@ -# -*- coding: utf-8 -*- -# Copyright (C) 2018-2024 Intel Corporation -# SPDX-License-Identifier: Apache-2.0 - -from openvino.utils.node_factory import NodeFactory diff --git a/src/bindings/python/src/openvino/utils/reduction.py b/src/bindings/python/src/openvino/runtime/utils/reduction.py similarity index 95% rename from src/bindings/python/src/openvino/utils/reduction.py rename to src/bindings/python/src/openvino/runtime/utils/reduction.py index e6be6d0ac9a104..71d0af8de7376e 100644 --- a/src/bindings/python/src/openvino/utils/reduction.py +++ b/src/bindings/python/src/openvino/runtime/utils/reduction.py @@ -4,7 +4,7 @@ from typing import Iterable, Optional -from openvino import Node +from openvino.runtime import Node def get_reduction_axes(node: Node, reduction_axes: Optional[Iterable[int]]) -> Iterable[int]: diff --git a/src/bindings/python/src/openvino/runtime/utils/reduction/__init__.py b/src/bindings/python/src/openvino/runtime/utils/reduction/__init__.py deleted file mode 100644 index a2fbff9e793dca..00000000000000 --- a/src/bindings/python/src/openvino/runtime/utils/reduction/__init__.py +++ /dev/null @@ -1,5 +0,0 @@ -# -*- coding: utf-8 -*- -# Copyright (C) 2018-2024 Intel Corporation -# SPDX-License-Identifier: Apache-2.0 - -from openvino.utils.reduction import get_reduction_axes diff --git a/src/bindings/python/src/openvino/utils/types.py b/src/bindings/python/src/openvino/runtime/utils/types.py similarity index 97% rename from src/bindings/python/src/openvino/utils/types.py rename to src/bindings/python/src/openvino/runtime/utils/types.py index b3543739741d94..52f1faf8e1e839 100644 --- a/src/bindings/python/src/openvino/utils/types.py +++ b/src/bindings/python/src/openvino/runtime/utils/types.py @@ -9,9 +9,9 @@ import numpy as np -from openvino.exceptions import OVTypeError -from openvino import Node, Shape, Output, Type -from openvino.op import Constant +from openvino.runtime.exceptions import OVTypeError +from openvino.runtime import Node, Shape, Output, Type +from openvino.runtime.op import Constant log = logging.getLogger(__name__) diff --git a/src/bindings/python/src/openvino/runtime/utils/types/__init__.py b/src/bindings/python/src/openvino/runtime/utils/types/__init__.py deleted file mode 100644 index 4f88d609988e8d..00000000000000 --- a/src/bindings/python/src/openvino/runtime/utils/types/__init__.py +++ /dev/null @@ -1,21 +0,0 @@ -# -*- coding: utf-8 -*- -# Copyright (C) 2018-2024 Intel Corporation -# SPDX-License-Identifier: Apache-2.0 - -from openvino.utils.types import TensorShape -from openvino.utils.types import NumericData -from openvino.utils.types import NumericType -from openvino.utils.types import ScalarData -from openvino.utils.types import NodeInput - -from openvino.utils.types import openvino_to_numpy_types_map -from openvino.utils.types import openvino_to_numpy_types_str_map -from openvino.utils.types import get_element_type -from openvino.utils.types import get_element_type_str -from openvino.utils.types import get_dtype -from openvino.utils.types import get_numpy_ctype -from openvino.utils.types import get_ndarray -from openvino.utils.types import get_shape -from openvino.utils.types import make_constant_node -from openvino.utils.types import as_node -from openvino.utils.types import as_nodes diff --git a/src/bindings/python/src/openvino/package_utils.py b/src/bindings/python/src/openvino/utils.py similarity index 97% rename from src/bindings/python/src/openvino/package_utils.py rename to src/bindings/python/src/openvino/utils.py index 6aa3f3ed39b556..9890ae9b3e6460 100644 --- a/src/bindings/python/src/openvino/package_utils.py +++ b/src/bindings/python/src/openvino/utils.py @@ -21,9 +21,9 @@ def _add_openvino_libs_to_search_path() -> None: if os.path.isdir(os.path.join(os.path.dirname(__file__), "libs")): # looking for the libs in the pip installation path. openvino_libs.append(os.path.join(os.path.dirname(__file__), "libs")) - elif os.path.isdir(os.path.join(os.path.dirname(__file__), os.pardir, os.pardir, os.pardir, "Library", "bin")): + elif os.path.isdir(os.path.join(os.path.dirname(__file__), "..", "..", "..", "Library", "bin")): # looking for the libs in the conda installation path - openvino_libs.append(os.path.join(os.path.dirname(__file__), os.pardir, os.pardir, os.pardir, "Library", "bin")) + openvino_libs.append(os.path.join(os.path.dirname(__file__), "..", "..", "..", "Library", "bin")) else: # setupvars.bat script set all libs paths to OPENVINO_LIB_PATHS environment variable. openvino_libs_installer = os.getenv("OPENVINO_LIB_PATHS") diff --git a/src/bindings/python/src/openvino/utils/__init__.py b/src/bindings/python/src/openvino/utils/__init__.py deleted file mode 100644 index 2ccc79d20cce84..00000000000000 --- a/src/bindings/python/src/openvino/utils/__init__.py +++ /dev/null @@ -1,12 +0,0 @@ -# -*- coding: utf-8 -*- -# Copyright (C) 2018-2024 Intel Corporation -# SPDX-License-Identifier: Apache-2.0 - -"""Generic utilities. Factor related functions out to separate files.""" - -from openvino._pyopenvino.util import numpy_to_c, replace_node, replace_output_update_name - -from openvino.package_utils import get_cmake_path -from openvino.package_utils import deprecated -from openvino.package_utils import classproperty -from openvino.package_utils import deprecatedclassproperty diff --git a/src/bindings/python/src/openvino/utils/data_helpers/__init__.py b/src/bindings/python/src/openvino/utils/data_helpers/__init__.py deleted file mode 100644 index 282547dd9df79a..00000000000000 --- a/src/bindings/python/src/openvino/utils/data_helpers/__init__.py +++ /dev/null @@ -1,8 +0,0 @@ -# -*- coding: utf-8 -*- -# Copyright (C) 2018-2024 Intel Corporation -# SPDX-License-Identifier: Apache-2.0 - -from openvino.utils.data_helpers.data_dispatcher import _data_dispatch -from openvino.utils.data_helpers.wrappers import tensor_from_file -from openvino.utils.data_helpers.wrappers import _InferRequestWrapper -from openvino.utils.data_helpers.wrappers import OVDict diff --git a/src/common/transformations/include/transformations/op_conversions/fake_convert_decomposition.hpp b/src/common/transformations/include/transformations/op_conversions/fake_convert_decomposition.hpp new file mode 100644 index 00000000000000..e149152b2bcf6d --- /dev/null +++ b/src/common/transformations/include/transformations/op_conversions/fake_convert_decomposition.hpp @@ -0,0 +1,32 @@ +// Copyright (C) 2024 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#pragma once + +#include "openvino/pass/matcher_pass.hpp" +#include "transformations_visibility.hpp" + +namespace ov { +namespace pass { + +class TRANSFORMATIONS_API FakeConvertDecomposition; + +} // namespace pass +} // namespace ov + +/** + * @ingroup ov_transformation_common_api + * @brief FakeConvertDecomposition transformation decomposes FakeConvert layer. + * f8: f8e4m3, f8e5m2 + * downconvert: f32->f8, f16->f8, bf16->f8 + * upconvert: f8->f32, f8->f16, f8->bf16 + * output = (upconvert(downconvert(input * scale - shift)) + shift) / scale + * + */ + +class ov::pass::FakeConvertDecomposition : public ov::pass::MatcherPass { +public: + OPENVINO_MATCHER_PASS_RTTI("FakeConvertDecomposition"); + FakeConvertDecomposition(); +}; diff --git a/src/common/transformations/src/transformations/op_conversions/fake_convert_decomposition.cpp b/src/common/transformations/src/transformations/op_conversions/fake_convert_decomposition.cpp new file mode 100644 index 00000000000000..7f0a44df6a151d --- /dev/null +++ b/src/common/transformations/src/transformations/op_conversions/fake_convert_decomposition.cpp @@ -0,0 +1,76 @@ +// Copyright (C) 2024 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "transformations/op_conversions/fake_convert_decomposition.hpp" + +#include "itt.hpp" +#include "openvino/core/rt_info.hpp" +#include "openvino/op/add.hpp" +#include "openvino/op/constant.hpp" +#include "openvino/op/convert.hpp" +#include "openvino/op/divide.hpp" +#include "openvino/op/fake_convert.hpp" +#include "openvino/op/multiply.hpp" +#include "openvino/op/subtract.hpp" +#include "openvino/pass/pattern/op/wrap_type.hpp" + +ov::pass::FakeConvertDecomposition::FakeConvertDecomposition() { + MATCHER_SCOPE(FakeConvertDecomposition); + auto data = pattern::any_input(); + + auto fake_convert = ov::pass::pattern::wrap_type(); + + matcher_pass_callback callback = [OV_CAPTURE_CPY_AND_THIS](ov::pass::pattern::Matcher& m) { + auto& pattern_to_output = m.get_pattern_value_map(); + const auto fake_convert_node = + ov::as_type_ptr(pattern_to_output.at(fake_convert).get_node_shared_ptr()); + + if (fake_convert_node == nullptr || transformation_callback(fake_convert_node)) { + return false; + } + + Output data{fake_convert_node->input_value(0)}; + const Output input_scale{fake_convert_node->input_value(1)}; + auto input_type = data.get_element_type(); + + ov::pass::NodeRegistry decomp_ops; + if (input_type != input_scale.get_element_type()) { + input_type = input_scale.get_element_type(); + data = std::make_shared(data, input_type); + data = decomp_ops.add(data.get_node_shared_ptr()); + } + + std::shared_ptr result; + const auto scale = decomp_ops.make(data, input_scale); + if (fake_convert_node->get_input_size() == 2) { + const auto downconvert = + decomp_ops.make(scale, fake_convert_node->get_destination_element_type()); + const auto upconvert = decomp_ops.make(downconvert, input_type); + + result = decomp_ops.make(upconvert, input_scale); + } else { + const Output input_shift{fake_convert_node->input_value(2)}; + const auto shift = decomp_ops.make(scale, input_shift); + + const auto downconvert = + decomp_ops.make(shift, fake_convert_node->get_destination_element_type()); + const auto upconvert = decomp_ops.make(downconvert, input_type); + + const auto deshift = decomp_ops.make(upconvert, input_shift); + result = decomp_ops.make(deshift, input_scale); + } + + if (result->get_output_element_type(0) != fake_convert_node->get_output_element_type(0)) { + result = decomp_ops.make(result, fake_convert_node->get_output_element_type(0)); + } + + result->set_friendly_name(m.get_match_root()->get_friendly_name()); + ov::copy_runtime_info(fake_convert_node, decomp_ops.get()); + ov::replace_node(m.get_match_root(), result); + return true; + }; + + auto m = std::make_shared(fake_convert, matcher_name); + register_matcher(m, callback); +} diff --git a/src/common/transformations/src/transformations/sdpa_to_paged_attention/position_ids_replacer.cpp b/src/common/transformations/src/transformations/sdpa_to_paged_attention/position_ids_replacer.cpp index 1cc9be37606950..397746c75bb84d 100644 --- a/src/common/transformations/src/transformations/sdpa_to_paged_attention/position_ids_replacer.cpp +++ b/src/common/transformations/src/transformations/sdpa_to_paged_attention/position_ids_replacer.cpp @@ -61,16 +61,19 @@ ov::pass::PositionIDsReplacerQwen::PositionIDsReplacerQwen(const Output& p auto p_opt_convert = optional(p_max_context_len); auto p_opt_reshape = optional({p_opt_convert, any_input()}); - // current seg len - auto p_input_ids = wrap_type(); - auto p_unsqueeze = wrap_type({p_input_ids, _const()}); - auto p_shape_of = wrap_type({p_unsqueeze}); + // current seq len: + // it might be present in 2 different ways: + // input_ids -> unsqueeze -> reshape -> convert -> shape_of -> gather + // QKV -> variadic_split(Q or K) -> rope Q/K -> shape_of -> gather + // Probably we can use the symbols to re-use one of these ways. + // Currently, "any_input" is used to detect the both places. + auto p_shape_of = wrap_type({any_input()}); auto p_current_len = wrap_type({p_shape_of, _const(), _const()}); - auto p_rotary_emb_sincos = wrap_type(); auto p_neg_const = wrap_type(); auto p_neg_mul = wrap_type({p_current_len, p_neg_const}); // the rotary_emb_cos/rotary_emb_sin are sliced by the total length [1,..4096,1,128] + auto p_rotary_emb_sincos = wrap_type(); auto p_slice_1 = wrap_type({p_rotary_emb_sincos, _const(), p_opt_reshape, _const(), _const()}); auto p_slice_2 = wrap_type({p_slice_1, p_neg_mul, _const(), _const(), _const()}); diff --git a/src/common/transformations/tests/op_conversions/fake_convert_decomposition_test.cpp b/src/common/transformations/tests/op_conversions/fake_convert_decomposition_test.cpp new file mode 100644 index 00000000000000..33b167ace11e24 --- /dev/null +++ b/src/common/transformations/tests/op_conversions/fake_convert_decomposition_test.cpp @@ -0,0 +1,149 @@ +// Copyright (C) 2024 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "transformations/op_conversions/fake_convert_decomposition.hpp" + +#include + +#include "common_test_utils/common_utils.hpp" +#include "common_test_utils/ov_test_utils.hpp" +#include "openvino/opsets/opset1.hpp" +#include "openvino/opsets/opset13.hpp" + +using namespace ov; + +using FakeConvertDecompositionParams = std::tuple; // default shift + +class FakeConvertDecompositionTest : public ov::test::TestsCommon, + public ::testing::WithParamInterface { +public: + static std::string getTestCaseName(::testing::TestParamInfo obj) { + FakeConvertDecompositionParams params = obj.param; + + Shape data_shape, scale_shape, shift_shape; + element::Type_t data_prec, dst_prec; + bool default_shift; + std::tie(data_shape, scale_shape, shift_shape, data_prec, dst_prec, default_shift) = params; + + std::ostringstream result; + result << "dataShape=" << ov::test::utils::vec2str(data_shape) << "_"; + result << "scaleShape=" << ov::test::utils::vec2str(scale_shape) << "_"; + result << "shiftShape=" << ov::test::utils::vec2str(shift_shape) << "_"; + result << "dataPrecision=" << element::Type(data_prec) << "_"; + result << "destinationPrecision=" << element::Type(dst_prec) << "_"; + if (default_shift) + result << "defaultShift=true"; + else + result << "defaultShift=false"; + return result.str(); + } +}; + +TEST_P(FakeConvertDecompositionTest, CompareFunctions) { + FakeConvertDecompositionParams params = this->GetParam(); + + Shape data_shape, scale_shape, shift_shape; + element::Type_t data_prec, dst_prec; + bool default_shift; + std::tie(data_shape, scale_shape, shift_shape, data_prec, dst_prec, default_shift) = params; + + std::shared_ptr model(nullptr); + { + const auto data = std::make_shared(data_prec, PartialShape(data_shape)); + const auto scale = std::make_shared(data_prec, scale_shape); + const auto shift = std::make_shared(data_prec, shift_shape); + + const auto fake_convert = default_shift ? std::make_shared(data, scale, dst_prec) + : std::make_shared(data, scale, shift, dst_prec); + model = std::make_shared(NodeVector{fake_convert}, ParameterVector{data}); + + pass::Manager manager; + manager.register_pass(); + manager.register_pass(); + manager.run_passes(model); + + OV_ASSERT_NO_THROW(check_rt_info(model)); + } + + std::shared_ptr model_ref(nullptr); + { + const auto input_data = std::make_shared(data_prec, PartialShape(data_shape)); + const auto input_scale = std::make_shared(data_prec, scale_shape); + const auto input_shift = std::make_shared(data_prec, shift_shape); + ParameterVector params; + params.push_back(input_data); + std::shared_ptr data = input_data; + + std::shared_ptr result; + const auto scale = std::make_shared(data, input_scale); + if (default_shift) { + const auto downconvert = std::make_shared(scale, dst_prec); + const auto upconvert = std::make_shared(downconvert, data_prec); + + result = std::make_shared(upconvert, input_scale); + } else { + const auto shift = std::make_shared(scale, input_shift); + + const auto downconvert = std::make_shared(shift, dst_prec); + const auto upconvert = std::make_shared(downconvert, data_prec); + + const auto deshift = std::make_shared(upconvert, input_shift); + result = std::make_shared(deshift, input_scale); + } + + model_ref = std::make_shared(NodeVector{result}, params); + } + + const auto res = compare_functions(model, model_ref); + ASSERT_TRUE(res.first) << res.second; +} + +const std::vector data_precisions = {element::Type_t::f32, + element::Type_t::f16, + element::Type_t::bf16}; + +const std::vector destination_precisions = {element::Type_t::f8e4m3, element::Type_t::f8e5m2}; + +const std::vector default_shift = {true, false}; + +const auto simple_fake_convert_params = ::testing::Combine(::testing::Values(Shape{2, 3, 4, 5}), + ::testing::Values(Shape{1}), + ::testing::Values(Shape{1}), + ::testing::ValuesIn(data_precisions), + ::testing::ValuesIn(destination_precisions), + ::testing::ValuesIn(default_shift)); + +const auto broadcast_fake_convert_params = ::testing::Combine(::testing::Values(Shape{2, 3, 4, 5}), + ::testing::Values(Shape{2, 3, 1, 1}), + ::testing::Values(Shape{2, 3, 1, 1}), + ::testing::ValuesIn(data_precisions), + ::testing::ValuesIn(destination_precisions), + ::testing::ValuesIn(default_shift)); + +const auto elementwise_fake_convert_params = ::testing::Combine(::testing::Values(Shape{2, 3, 4, 5}), + ::testing::Values(Shape{2, 3, 4, 5}), + ::testing::Values(Shape{2, 3, 4, 5}), + ::testing::ValuesIn(data_precisions), + ::testing::ValuesIn(destination_precisions), + ::testing::ValuesIn(default_shift)); + +INSTANTIATE_TEST_SUITE_P(SimpleFakeConvert_Decomposition, + FakeConvertDecompositionTest, + simple_fake_convert_params, + FakeConvertDecompositionTest::getTestCaseName); + +INSTANTIATE_TEST_SUITE_P(BroadcastFakeConvert_Decomposition, + FakeConvertDecompositionTest, + broadcast_fake_convert_params, + FakeConvertDecompositionTest::getTestCaseName); + +INSTANTIATE_TEST_SUITE_P(ElementwiseFakeConvert_Decomposition, + FakeConvertDecompositionTest, + elementwise_fake_convert_params, + FakeConvertDecompositionTest::getTestCaseName); diff --git a/src/frontends/onnx/tests/__init__.py b/src/frontends/onnx/tests/__init__.py index ef8cebfa361e3f..fdf1295dfd1dbe 100644 --- a/src/frontends/onnx/tests/__init__.py +++ b/src/frontends/onnx/tests/__init__.py @@ -147,7 +147,7 @@ def xfail_test(reason="Mark the test as expected to fail", strict=True): skip_dynamic_model = pytest.mark.skip(reason="CPU plug-in can't load a model with dynamic output shapes via legacy API") # ONNX 1.14 -xfail_issue_119896 = xfail_test(reason="Unsupported element type: FLOAT8") +xfail_issue_119896 = xfail_test(reason="Unsupported element type: FLOAT8", strict=False) xfail_issue_119900 = xfail_test(reason="While validating ONNX node '': " "half_pixel_symmetric - this type of coordinate transformation mode " "is not supported. Choose one of the following modes: " diff --git a/src/inference/src/os/lin/lin_system_conf.cpp b/src/inference/src/os/lin/lin_system_conf.cpp index 64da4cb0ac836a..29c8bfddbd1ca4 100644 --- a/src/inference/src/os/lin/lin_system_conf.cpp +++ b/src/inference/src/os/lin/lin_system_conf.cpp @@ -219,14 +219,16 @@ CPU::CPU() { } else if (valid_cpu_mapping_table.size() == (unsigned)_processors) { return 0; } else { - std::lock_guard lock{_cpu_mutex}; _processors = valid_cpu_mapping_table.size(); _cpu_mapping_table.swap(valid_cpu_mapping_table); - update_valid_processor_linux(std::move(phy_core_list), - _numa_nodes, - _cores, - _proc_type_table, - _cpu_mapping_table); + { + std::lock_guard lock{_cpu_mutex}; + update_valid_processor_linux(std::move(phy_core_list), + _numa_nodes, + _cores, + _proc_type_table, + _cpu_mapping_table); + } return 0; } }; @@ -235,7 +237,7 @@ CPU::CPU() { if (!get_info_linux(cache_info_mode)) { parse_cache_info_linux(system_info_table, - node_info_table, + std::move(node_info_table), _processors, _numa_nodes, _sockets, @@ -249,7 +251,7 @@ CPU::CPU() { (_proc_type_table[0][ALL_PROC] != _proc_type_table[0][EFFICIENT_CORE_PROC]))) { if (!get_info_linux(freq_info_mode)) { parse_freq_info_linux(system_info_table, - node_info_table, + std::move(node_info_table), _processors, _numa_nodes, _sockets, diff --git a/src/plugins/intel_cpu/src/dnnl_extension_utils.cpp b/src/plugins/intel_cpu/src/dnnl_extension_utils.cpp index 457f8368f734dd..1c5598b6d55e26 100644 --- a/src/plugins/intel_cpu/src/dnnl_extension_utils.cpp +++ b/src/plugins/intel_cpu/src/dnnl_extension_utils.cpp @@ -36,6 +36,8 @@ uint8_t DnnlExtensionUtils::sizeOfDataType(dnnl::memory::data_type dataType) { case dnnl::memory::data_type::s4: case dnnl::memory::data_type::u4: case dnnl::memory::data_type::f8_e8m0: + case dnnl::memory::data_type::f8_e4m3: + case dnnl::memory::data_type::f8_e5m2: case dnnl::memory::data_type::f4_e2m1: return 1; case dnnl::memory::data_type::undef: @@ -70,6 +72,10 @@ dnnl::memory::data_type DnnlExtensionUtils::ElementTypeToDataType(const ov::elem return memory::data_type::u4; case ov::element::f8e8m0: return memory::data_type::f8_e8m0; + case ov::element::f8e4m3: + return memory::data_type::f8_e4m3; + case ov::element::f8e5m2: + return memory::data_type::f8_e5m2; case ov::element::f4e2m1: return memory::data_type::f4_e2m1; case ov::element::undefined: @@ -106,6 +112,10 @@ ov::element::Type DnnlExtensionUtils::DataTypeToElementType(const dnnl::memory:: return ov::element::u4; case memory::data_type::f8_e8m0: return ov::element::f8e8m0; + case memory::data_type::f8_e4m3: + return ov::element::f8e4m3; + case memory::data_type::f8_e5m2: + return ov::element::f8e5m2; case memory::data_type::f4_e2m1: return ov::element::f4e2m1; case memory::data_type::undef: diff --git a/src/plugins/intel_cpu/src/nodes/common/cpu_convert.cpp b/src/plugins/intel_cpu/src/nodes/common/cpu_convert.cpp index 0c8cddd905dc2e..f6aabe376d6eec 100644 --- a/src/plugins/intel_cpu/src/nodes/common/cpu_convert.cpp +++ b/src/plugins/intel_cpu/src/nodes/common/cpu_convert.cpp @@ -9,6 +9,7 @@ #include "utils/bfloat16.hpp" #if defined(OPENVINO_ARCH_X86_64) +# include "cpu/x64/jit_avx512_core_fp8cvt.hpp" # include "nodes/kernels/x64/jit_kernel.hpp" #else # include "cpu_memory.h" @@ -27,6 +28,18 @@ using namespace dnnl::impl::utils; using namespace dnnl::impl::cpu::x64; using namespace Xbyak; +enum f8_type { none, f8e4m3, f8e5m2 }; + +template +f8_type get_f8_type() { + if (std::is_same::value || std::is_same::value) { + return f8_type::f8e4m3; + } else if (std::is_same::value || std::is_same::value) { + return f8_type::f8e5m2; + } + return f8_type::none; +} + template void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst); @@ -50,12 +63,14 @@ void convert_vec(jit_generator& gen, const RegExp& src, cons gen.movdqu(gen.xword[dst], f16vec); } +template class jit_convert_array : public jit_kernel { DECLARE_CPU_JIT_AUX_FUNCTIONS(jit_convert_array) void generate() override { - constexpr size_t vlen = 8u; - constexpr size_t vlen_log2 = 3; + bool is_fp8 = f8_e4m3_emu_ || f8_e5m2_emu_; + size_t vlen = is_fp8 ? 16u : 8u; + size_t vlen_log2 = is_fp8 ? 4 : 3; preamble(); @@ -84,17 +99,24 @@ class jit_convert_array : public jit_kernel { auto tail_size = var(); tail_size = size; - tail_size <<= static_cast(std::logb(_src_size)) - 1; - copy(tmp.pointer(), src, tail_size); + tail_size <<= static_cast(std::logb(_src_size)); + copy(tmp.pointer(), src, tail_size); _convert_vec(*this, tmp.pointer(), tmp.pointer()); tail_size = size; - tail_size <<= static_cast(std::logb(_dst_size)) - 1; - copy(dst, tmp.pointer(), tail_size); + tail_size <<= static_cast(std::logb(_dst_size)); + copy(dst, tmp.pointer(), tail_size); }); postamble(); + + if (f8_e4m3_emu_) + f8_e4m3_emu_->prepare_table(); + if (f8_e5m2_emu_) + f8_e5m2_emu_->prepare_table(); + if (uni_vcvtneps2bf16_) + uni_vcvtneps2bf16_->emit_data(); } public: @@ -108,16 +130,37 @@ class jit_convert_array : public jit_kernel { typedef void (*convert_vec_t)(jit_generator&, const RegExp&, const RegExp&); - jit_convert_array(convert_vec_t convert_vec, size_t src_size, size_t dst_size) + jit_convert_array(convert_vec_t convert_vec) : jit_kernel(jit_name()), _convert_vec(convert_vec), - _src_size(src_size), - _dst_size(dst_size) {} + _src_size(sizeof(src_t)), + _dst_size(sizeof(dst_t)) { + const auto type = get_f8_type(); + if (type == f8_type::f8e4m3) { + f8_e4m3_emu_ = std::make_shared(this, + fp8_emu_reserv_1_, + fp8_emu_reserv_2_, + fp8_emu_reserv_3_, + fp8_emu_reserv_4_, + fp8_emu_reserv_5_, + fp8_emu_scratch_); + } else if (type == f8_type::f8e5m2) { + f8_e5m2_emu_ = std::make_shared(this, + fp8_emu_reserv_1_, + fp8_emu_reserv_2_, + fp8_emu_reserv_3_, + fp8_emu_kmask_aux_, + fp8_emu_scratch_); + } + const bool is_dst_bf16 = std::is_same::value; + if (is_dst_bf16 && mayiuse(cpu_isa_t::avx512_core)) { + uni_vcvtneps2bf16_ = std::make_shared(this, cpu_isa_t::avx512_core); + } + } - template static fn_t get() { if (mayiuse(cpu_isa_t::avx2) && dnnl::impl::cpu::x64::cpu().has(Xbyak::util::Cpu::tF16C)) { - static jit_convert_array converter(convert_vec, sizeof(src_t), sizeof(dst_t)); + static jit_convert_array converter(convert_vec); auto& generator = static_cast(converter); generator.create_kernel(); return (fn_t)generator.jit_ker(); @@ -125,16 +168,192 @@ class jit_convert_array : public jit_kernel { return nullptr; } + std::shared_ptr get_f8_e4m3_emu() const { + return f8_e4m3_emu_; + } + + std::shared_ptr get_f8_e5m2_emu() const { + return f8_e5m2_emu_; + } + + std::shared_ptr get_uni_vcvtneps2bf16() const { + return uni_vcvtneps2bf16_; + } + private: convert_vec_t _convert_vec; size_t _src_size; size_t _dst_size; + + std::shared_ptr f8_e4m3_emu_; + std::shared_ptr f8_e5m2_emu_; + std::shared_ptr uni_vcvtneps2bf16_; + + const Reg64 fp8_emu_scratch_ = rax; + const Zmm fp8_emu_reserv_1_ = Zmm(9); + const Zmm fp8_emu_reserv_2_ = Zmm(10); + const Zmm fp8_emu_reserv_3_ = Zmm(11); + const Zmm fp8_emu_reserv_4_ = Zmm(12); + const Zmm fp8_emu_reserv_5_ = Zmm(13); + const Opmask fp8_emu_kmask_aux_ = Opmask(1); }; +template <> +void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst) { + auto const& f8vec = gen.xmm3; + auto const& f32vec = gen.zmm4; + + auto& cvt = dynamic_cast&>(gen); + + gen.vmovups(f32vec, gen.zword[src]); + cvt.get_f8_e4m3_emu()->vcvt_f32_to_f8(f8vec, f32vec); + gen.vmovdqu(gen.xword[dst], f8vec); +} + +template <> +void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst) { + auto const& f8vec = gen.xmm3; + auto const& f32vec = gen.zmm4; + + auto& cvt = dynamic_cast&>(gen); + + gen.vmovdqu(f8vec, gen.xword[src]); + cvt.get_f8_e4m3_emu()->vcvt_f8_to_f32(f32vec, f8vec); + gen.vmovups(gen.zword[dst], f32vec); +} + +template <> +void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst) { + auto const& f8vec = gen.xmm3; + auto const& f16vec = gen.ymm4; + + auto& cvt = dynamic_cast&>(gen); + + gen.vmovdqu(f16vec, gen.yword[src]); + cvt.get_f8_e4m3_emu()->vcvt_f16_to_f8(f8vec, f16vec); + gen.vmovdqu(gen.xword[dst], f8vec); +} + +template <> +void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst) { + auto const& f8vec = gen.xmm3; + auto const& f16vec = gen.ymm4; + + auto& cvt = dynamic_cast&>(gen); + + gen.vmovdqu(f8vec, gen.xword[src]); + cvt.get_f8_e4m3_emu()->vcvt_f8_to_f16(f16vec, f8vec); + gen.vmovdqu(gen.yword[dst], f16vec); +} + +template <> +void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst) { + auto const& f8vec = gen.xmm3; + auto const& f16vec = gen.zmm4; + + auto& cvt = dynamic_cast&>(gen); + + gen.vpmovzxwd(f16vec, gen.yword[src]); + gen.vpslld(f16vec, f16vec, 16); + cvt.get_f8_e4m3_emu()->vcvt_f32_to_f8(f8vec, f16vec); + gen.vmovdqu(gen.xword[dst], f8vec); +} + +template <> +void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst) { + auto const& f8vec = gen.xmm3; + auto const& f16vec = gen.ymm4; + auto const& f32vec = gen.zmm4; + + auto& cvt = dynamic_cast&>(gen); + + gen.vmovdqu(f8vec, gen.xword[src]); + cvt.get_f8_e4m3_emu()->vcvt_f8_to_f32(f32vec, f8vec); + cvt.get_uni_vcvtneps2bf16()->emit_code({static_cast(f32vec.getIdx())}, + {static_cast(f16vec.getIdx())}); + gen.vmovdqu(gen.yword[dst], f16vec); +} + +template <> +void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst) { + auto const& f8vec = gen.xmm3; + auto const& f32vec = gen.zmm4; + + auto& cvt = dynamic_cast&>(gen); + + gen.vmovups(f32vec, gen.zword[src]); + cvt.get_f8_e5m2_emu()->vcvt_f32_to_f8(f8vec, f32vec); + gen.vmovdqu(gen.xword[dst], f8vec); +} + +template <> +void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst) { + auto const& f8vec = gen.xmm3; + auto const& f32vec = gen.zmm4; + + auto& cvt = dynamic_cast&>(gen); + + gen.vmovdqu(f8vec, gen.xword[src]); + cvt.get_f8_e5m2_emu()->vcvt_f8_to_f32(f32vec, f8vec); + gen.vmovups(gen.zword[dst], f32vec); +} + +template <> +void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst) { + auto const& f8vec = gen.xmm3; + auto const& f16vec = gen.ymm4; + + auto& cvt = dynamic_cast&>(gen); + + gen.vmovdqu(f16vec, gen.yword[src]); + cvt.get_f8_e5m2_emu()->vcvt_f16_to_f8(f8vec, f16vec); + gen.vmovdqu(gen.xword[dst], f8vec); +} + +template <> +void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst) { + auto const& f8vec = gen.xmm3; + auto const& f16vec = gen.ymm4; + + auto& cvt = dynamic_cast&>(gen); + + gen.vmovdqu(f8vec, gen.xword[src]); + cvt.get_f8_e5m2_emu()->vcvt_f8_to_f16(f16vec, f8vec); + gen.vmovdqu(gen.yword[dst], f16vec); +} + +template <> +void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst) { + auto const& f8vec = gen.xmm3; + auto const& f16vec = gen.zmm4; + + auto& cvt = dynamic_cast&>(gen); + + gen.vpmovzxwd(f16vec, gen.yword[src]); + gen.vpslld(f16vec, f16vec, 16); + cvt.get_f8_e5m2_emu()->vcvt_f32_to_f8(f8vec, f16vec); + gen.vmovdqu(gen.xword[dst], f8vec); +} + +template <> +void convert_vec(jit_generator& gen, const RegExp& src, const RegExp& dst) { + auto const& f8vec = gen.xmm3; + auto const& f16vec = gen.ymm4; + auto const& f32vec = gen.zmm4; + + auto& cvt = dynamic_cast&>(gen); + + gen.vmovdqu(f8vec, gen.xword[src]); + cvt.get_f8_e5m2_emu()->vcvt_f8_to_f32(f32vec, f8vec); + cvt.get_uni_vcvtneps2bf16()->emit_code({static_cast(f32vec.getIdx())}, + {static_cast(f16vec.getIdx())}); + gen.vmovdqu(gen.yword[dst], f16vec); +} + template void jit_convert(const TI* arg, TO* out, size_t count) { - using jit_impl = jit_convert_array; - static auto converter = jit_impl::get(); + using jit_impl = jit_convert_array; + static auto converter = jit_impl::get(); if (converter) { typename jit_impl::args_t args = {arg, out, count}; @@ -185,6 +404,12 @@ const std::tuple& Range::fit(const ov::element::Type& prec) { if (prec.is_real()) { double lbound, ubound; switch (prec) { + case ov::element::f8e4m3: + lbound = static_cast(std::numeric_limits::lowest()); + ubound = static_cast(std::numeric_limits::max()); + case ov::element::f8e5m2: + lbound = static_cast(std::numeric_limits::lowest()); + ubound = static_cast(std::numeric_limits::max()); case ov::element::bf16: lbound = static_cast(std::numeric_limits::lowest()); ubound = static_cast(std::numeric_limits::max()); @@ -293,6 +518,18 @@ struct ConvertPrecision> { src_t lbound, ubound; std::tie(lbound, ubound) = ctx.range(); + // Align with the behavior of ngraph ref and jit implementation. Conversion from f8e4m3-inf + // to float should output float-inf instead of f8e4m3-max. Proper handling of special values + // (nan, inf, overflow) has already been assured by the conversion process. + if (std::is_same::value || std::is_same::value || + std::is_same::value || std::is_same::value) { + parallel_for(ctx.size, [&](size_t i) { + dst[i] = static_cast(src[i]); + }); + ctx.converted = true; + return; + } + if (std::is_integral::value || ctx.interimPrc.is_real() || std::is_integral::value) { parallel_for(ctx.size, [&](size_t i) { dst[i] = static_cast(std::max(std::min(src[i], ubound), lbound)); @@ -492,6 +729,12 @@ struct ConvertPrecision> { PrecisionInfo::value_type, \ PrecisionInfo::value_type) +#define INTEL_CPU_CVT_FP8_LIST \ + INTEL_CPU_CVT(f32, f8e4m3), INTEL_CPU_CVT(f16, f8e4m3), INTEL_CPU_CVT(bf16, f8e4m3), INTEL_CPU_CVT(f8e4m3, f32), \ + INTEL_CPU_CVT(f8e4m3, f16), INTEL_CPU_CVT(f8e4m3, bf16), INTEL_CPU_CVT(f32, f8e5m2), \ + INTEL_CPU_CVT(f16, f8e5m2), INTEL_CPU_CVT(bf16, f8e5m2), INTEL_CPU_CVT(f8e5m2, f32), \ + INTEL_CPU_CVT(f8e5m2, f16), INTEL_CPU_CVT(f8e5m2, bf16) + #define INTEL_CPU_CVT_LIST \ INTEL_CPU_CVT(u8, i8), INTEL_CPU_CVT(u8, u16), INTEL_CPU_CVT(u8, i16), INTEL_CPU_CVT(u8, u32), \ INTEL_CPU_CVT(u8, i32), INTEL_CPU_CVT(u8, u64), INTEL_CPU_CVT(u8, i64), INTEL_CPU_CVT(u8, f32), \ @@ -535,7 +778,8 @@ struct ConvertPrecision> { INTEL_CPU_CVT(boolean, f16), INTEL_CPU_CVT(boolean, bf16), INTEL_CPU_CVT(boolean, f64), INTEL_CPU_CVT(u8, u8), \ INTEL_CPU_CVT(i8, i8), INTEL_CPU_CVT(u16, u16), INTEL_CPU_CVT(i16, i16), INTEL_CPU_CVT(u32, u32), \ INTEL_CPU_CVT(i32, i32), INTEL_CPU_CVT(u64, u64), INTEL_CPU_CVT(i64, i64), INTEL_CPU_CVT(f32, f32), \ - INTEL_CPU_CVT(f16, f16), INTEL_CPU_CVT(bf16, bf16), INTEL_CPU_CVT(f64, f64), INTEL_CPU_CVT(boolean, boolean) + INTEL_CPU_CVT(f16, f16), INTEL_CPU_CVT(bf16, bf16), INTEL_CPU_CVT(f64, f64), INTEL_CPU_CVT(boolean, boolean), \ + INTEL_CPU_CVT_FP8_LIST #define INTEL_CPU_CVT_FROM_BIN_LIST \ INTEL_CPU_CVT(u1, f32), INTEL_CPU_CVT(u1, f16), INTEL_CPU_CVT(u1, bf16), INTEL_CPU_CVT(u1, f64), \ @@ -667,6 +911,35 @@ struct ConvertFromByteFPPrecision> { } }; +#if defined(OPENVINO_ARCH_X86_64) +struct ConvertFP8Context { + const void* srcPtr; + void* dstPtr; + size_t size; + bool converted; +}; + +template +struct ConvertFP8Precision; + +template +struct ConvertFP8Precision> { + void operator()(ConvertFP8Context& ctx) { + auto src = static_cast(ctx.srcPtr); + auto dst = static_cast(ctx.dstPtr); + constexpr size_t batch = 64; + const size_t iterations = ov::intel_cpu::div_up(ctx.size, batch); + parallel_for(iterations, [&](size_t i) { + const size_t offset = i * batch; + const size_t current_batch_size = std::min(ctx.size - offset, batch); + jit_convert(src + offset, dst + offset, current_batch_size); + }); + + ctx.converted = true; + } +}; +#endif + void cpu_convert(const void* srcPtr, void* dstPtr, ov::element::Type srcPrc, @@ -728,7 +1001,7 @@ void cpu_convert(const void* srcPtr, OV_SWITCH(intel_cpu, ConvertFrom4BitPrecision, ctx, std::tie(srcPrc, dstPrc), INTEL_CPU_CVT_FROM_4BIT_LIST); if (!ctx.converted) OPENVINO_THROW("cpu_convert can't convert from: ", srcPrc, " precision to: ", dstPrc); - } else if (srcPrc.bitwidth() == 8u && srcPrc.is_real()) { + } else if (srcPrc == ov::element::f8e8m0) { ConvertFromByteFPContext ctx{srcPrc, srcPtr, dstPtr, size, false}; OV_SWITCH(intel_cpu, ConvertFromByteFPPrecision, @@ -737,6 +1010,15 @@ void cpu_convert(const void* srcPtr, INTEL_CPU_CVT_FROM_BYTE_FP_LIST); if (!ctx.converted) OPENVINO_THROW("cpu_convert can't convert from: ", srcPrc, " precision to: ", dstPrc); +#if defined(OPENVINO_ARCH_X86_64) + } else if (dnnl::impl::cpu::x64::mayiuse(dnnl::impl::cpu::x64::avx512_core_fp16) && + (one_of(srcPrc, ov::element::f8e4m3, ov::element::f8e5m2) || + one_of(dstPrc, ov::element::f8e4m3, ov::element::f8e5m2))) { + ConvertFP8Context ctx{srcPtr, dstPtr, size, false}; + OV_SWITCH(intel_cpu, ConvertFP8Precision, ctx, std::tie(srcPrc, dstPrc), INTEL_CPU_CVT_FP8_LIST); + if (!ctx.converted) + OPENVINO_THROW("cpu_convert can't convert from: ", srcPrc, " precision to: ", dstPrc); +#endif } else { ConvertContext ctx{srcPtr, dstPtr, size, interimPrc, dstPrc, false}; OV_SWITCH(intel_cpu, ConvertPrecision, ctx, std::tie(srcPrc, dstPrc), INTEL_CPU_CVT_LIST); diff --git a/src/plugins/intel_cpu/src/plugin.cpp b/src/plugins/intel_cpu/src/plugin.cpp index db55c728df725e..b3c2aa0b298a5a 100644 --- a/src/plugins/intel_cpu/src/plugin.cpp +++ b/src/plugins/intel_cpu/src/plugin.cpp @@ -218,6 +218,8 @@ std::shared_ptr Plugin::compile_model(const std::shared_ptr< ov::element::Type_t::i4, ov::element::Type_t::u8, ov::element::Type_t::i8, + ov::element::Type_t::f8e4m3, + ov::element::Type_t::f8e5m2, ov::element::Type_t::u16, ov::element::Type_t::i16, ov::element::Type_t::u32, diff --git a/src/plugins/intel_cpu/src/transformations/transformation_pipeline.cpp b/src/plugins/intel_cpu/src/transformations/transformation_pipeline.cpp index fb9e0925bc89e2..4d7df9a335e98a 100644 --- a/src/plugins/intel_cpu/src/transformations/transformation_pipeline.cpp +++ b/src/plugins/intel_cpu/src/transformations/transformation_pipeline.cpp @@ -80,6 +80,7 @@ #include "transformations/op_conversions/detection_output_downgrade.hpp" #include "transformations/op_conversions/detection_output_upgrade.hpp" #include "transformations/op_conversions/eye_decomposition.hpp" +#include "transformations/op_conversions/fake_convert_decomposition.hpp" #include "transformations/op_conversions/fq_decomposition.hpp" #include "transformations/op_conversions/gelu7_downgrade.hpp" #include "transformations/op_conversions/group_normalization_decomposition.hpp" @@ -1293,6 +1294,7 @@ void Transformations::PostSnippets(void) { return node::FakeQuantize::isSupportedOperation(node, errMsg); }, ov::pass::FakeQuantizeDecomposition); + CPU_REGISTER_PASS_COMMON(postSnippetsManager, ov::pass::FakeConvertDecomposition); CPU_REGISTER_PASS_COMMON(postSnippetsManager, ov::pass::ConstantFolding); postSnippetsManager.run_passes(model); } diff --git a/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/classes/conversion.cpp b/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/classes/conversion.cpp index 4989fb3a0f04b7..a3c1f9ef7d3544 100644 --- a/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/classes/conversion.cpp +++ b/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/classes/conversion.cpp @@ -16,11 +16,45 @@ using namespace CPUTestUtils; namespace ov { namespace test { +static std::string special_value_to_string(const ov::test::SpecialValue& value) { + if (value == SpecialValue::none) { + return "none"; + } else if (value == SpecialValue::nan) { + return "nan"; + } else if (value == SpecialValue::inf) { + return "inf"; + } else if (value == SpecialValue::overflow) { + return "overflow"; + } + return "unknown"; +} + +template +static T set_special_value(T& value, const ov::test::SpecialValue& special_value) { + if (special_value == ov::test::SpecialValue::nan) { + value = NAN; + } else if (special_value == ov::test::SpecialValue::inf) { + value = INFINITY; + } else if (special_value == ov::test::SpecialValue::overflow) { + value = value + std::numeric_limits::max(); + } + return value; +} + +template +static void modify_value(ov::Tensor& tensor, const ov::test::SpecialValue& special_value) { + T* dataPtr = static_cast(tensor.data()); + for (size_t i = 0; i < tensor.get_size(); i++) { + set_special_value(dataPtr[i], special_value); + } +} + std::string ConvertCPULayerTest::getTestCaseName(testing::TestParamInfo obj) { InputShape inputShape; ov::element::Type inPrc, outPrc; + ov::test::SpecialValue special_value; CPUSpecificParams cpuParams; - std::tie(inputShape, inPrc, outPrc, cpuParams) = obj.param; + std::tie(inputShape, inPrc, outPrc, special_value, cpuParams) = obj.param; std::ostringstream result; @@ -30,6 +64,7 @@ std::string ConvertCPULayerTest::getTestCaseName(testing::TestParamInfo(inPrc, shape)); @@ -101,6 +146,31 @@ void ConvertCPULayerTest::SetUp() { function = makeNgraphFunction(inPrc, params, conversion, "ConversionCPU"); } +void ConvertCPULayerTest::generate_inputs(const std::vector& targetInputStaticShapes) { + inputs.clear(); + const auto& funcInputs = function->inputs(); + for (size_t i = 0; i < funcInputs.size(); ++i) { + const auto& funcInput = funcInputs[i]; + ov::Tensor tensor = + ov::test::utils::create_and_fill_tensor(funcInput.get_element_type(), targetInputStaticShapes[i]); + if (special_value != ov::test::SpecialValue::none) { + if (inPrc == ov::element::f32) { + modify_value(tensor, special_value); + } else if (inPrc == ov::element::f16) { + modify_value(tensor, special_value); + } else if (inPrc == ov::element::bf16) { + modify_value(tensor, special_value); + } else if (inPrc == ov::element::f8e4m3) { + modify_value(tensor, special_value); + } else if (inPrc == ov::element::f8e5m2) { + modify_value(tensor, special_value); + } + } + + inputs.insert({funcInput.get_node_shared_ptr(), tensor}); + } +} + void ConvertCPULayerTest::validate_out_prc() const { if (outPrc == ov::element::boolean) FAIL() << "ConvertCPULayerTest supports only non boolean output prc"; diff --git a/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/classes/conversion.hpp b/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/classes/conversion.hpp index a53f56f873151c..a4f4e0fc56c238 100644 --- a/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/classes/conversion.hpp +++ b/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/classes/conversion.hpp @@ -13,9 +13,12 @@ using namespace CPUTestUtils; namespace ov { namespace test { +enum SpecialValue { none, nan, inf, overflow }; + using convertLayerTestParamsSet = std::tuple; class ConvertCPULayerTest : public testing::WithParamInterface, @@ -25,9 +28,12 @@ class ConvertCPULayerTest : public testing::WithParamInterface& targetInputStaticShapes) override; virtual void validate_out_prc() const; ov::element::Type inPrc, outPrc; +private: + ov::test::SpecialValue special_value; }; class ConvertToBooleanCPULayerTest : public ConvertCPULayerTest { diff --git a/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/instances/arm/conversion.cpp b/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/instances/arm/conversion.cpp index 11e0440b2e3618..e5d87f5cb2f3dd 100644 --- a/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/instances/arm/conversion.cpp +++ b/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/instances/arm/conversion.cpp @@ -16,6 +16,7 @@ INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_7D_Dynamic, ConvertCPULayerTe ::testing::ValuesIn(inShapes_7D_dynamic()), ::testing::ValuesIn(precisions()), ::testing::ValuesIn(precisions()), + ::testing::Values(ov::test::SpecialValue::none), ::testing::Values(CPUSpecificParams({}, {}, {}, {}))), ConvertCPULayerTest::getTestCaseName); @@ -24,6 +25,7 @@ INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_7D_Static, ConvertCPULayerTes ::testing::ValuesIn(inShapes_7D_static()), ::testing::ValuesIn(precisions()), ::testing::ValuesIn(precisions()), + ::testing::Values(ov::test::SpecialValue::none), ::testing::Values(CPUSpecificParams({}, {}, {}, {}))), ConvertCPULayerTest::getTestCaseName); diff --git a/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/instances/common/conversion.cpp b/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/instances/common/conversion.cpp index 59ca1065bf78d9..8181304bf95e7d 100644 --- a/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/instances/common/conversion.cpp +++ b/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/instances/common/conversion.cpp @@ -31,6 +31,7 @@ INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_4D_Dynamic, ConvertCPULayerTe ::testing::ValuesIn(inShapes_4D_dynamic()), ::testing::ValuesIn(precisions()), ::testing::ValuesIn(precisions()), + ::testing::Values(ov::test::SpecialValue::none), ::testing::ValuesIn(memForm4D_dynamic)), ConvertCPULayerTest::getTestCaseName); @@ -39,6 +40,7 @@ INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_4bit_Dynamic, ConvertCPULayer ::testing::Combine(::testing::ValuesIn(inShapes_4D_dynamic()), ::testing::ValuesIn({ov::element::u4, ov::element::i4}), ::testing::ValuesIn({ov::element::f32, ov::element::bf16, ov::element::u8, ov::element::i8}), + ::testing::Values(ov::test::SpecialValue::none), ::testing::Values(CPUSpecificParams({nchw}, {nchw}, {}, {"ref"}))), ConvertCPULayerTest::getTestCaseName); @@ -52,9 +54,69 @@ INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_4D_Static, ConvertCPULayerTes ::testing::ValuesIn(inShapes_4D_static()), ::testing::ValuesIn(precisions()), ::testing::ValuesIn(precisions()), + ::testing::Values(ov::test::SpecialValue::none), ::testing::ValuesIn(memForm4D_static_common)), ConvertCPULayerTest::getTestCaseName); +const std::vector float_precisions = { + ov::element::f32, + ov::element::f16, + ov::element::bf16, +}; + +const std::vector f8_precisions = { + ov::element::f8e4m3, + ov::element::f8e5m2, +}; + +const std::vector specialValue = { + ov::test::SpecialValue::none, + ov::test::SpecialValue::nan, + ov::test::SpecialValue::inf, + ov::test::SpecialValue::overflow, +}; + +std::vector memForm4D_fp8 = { + CPUSpecificParams({nchw}, {nchw}, {}, expectedPrimitiveType()), + CPUSpecificParams({nhwc}, {nhwc}, {}, expectedPrimitiveType()), +}; + +INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_from_fp8_Static, ConvertCPULayerTest, + ::testing::Combine( + ::testing::ValuesIn(inShapes_4D_static()), + ::testing::ValuesIn(f8_precisions), + ::testing::ValuesIn(float_precisions), + ::testing::ValuesIn(specialValue), + ::testing::ValuesIn(memForm4D_fp8)), + ConvertCPULayerTest::getTestCaseName); + +INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_to_fp8_Static, ConvertCPULayerTest, + ::testing::Combine( + ::testing::ValuesIn(inShapes_4D_static()), + ::testing::ValuesIn(float_precisions), + ::testing::ValuesIn(f8_precisions), + ::testing::ValuesIn(specialValue), + ::testing::ValuesIn(memForm4D_fp8)), + ConvertCPULayerTest::getTestCaseName); + +INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_from_fp8_Dynamic, ConvertCPULayerTest, + ::testing::Combine( + ::testing::ValuesIn(inShapes_4D_dynamic()), + ::testing::ValuesIn(f8_precisions), + ::testing::ValuesIn(float_precisions), + ::testing::ValuesIn(specialValue), + ::testing::ValuesIn(memForm4D_fp8)), + ConvertCPULayerTest::getTestCaseName); + +INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_to_fp8_Dynamic, ConvertCPULayerTest, + ::testing::Combine( + ::testing::ValuesIn(inShapes_4D_dynamic()), + ::testing::ValuesIn(float_precisions), + ::testing::ValuesIn(f8_precisions), + ::testing::ValuesIn(specialValue), + ::testing::ValuesIn(memForm4D_fp8)), + ConvertCPULayerTest::getTestCaseName); + } // namespace Conversion } // namespace test } // namespace ov \ No newline at end of file diff --git a/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/instances/x64/conversion.cpp b/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/instances/x64/conversion.cpp index 9c34d6220d4b2d..ab1e06639c5a3e 100644 --- a/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/instances/x64/conversion.cpp +++ b/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/instances/x64/conversion.cpp @@ -23,6 +23,7 @@ INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_blocked_Dynamic, ConvertCPULa ::testing::ValuesIn(inShapes_4D_dynamic()), ::testing::ValuesIn(precisions()), ::testing::ValuesIn(precisions()), + ::testing::Values(ov::test::SpecialValue::none), ::testing::ValuesIn(memForm4D_dynamic)), ConvertCPULayerTest::getTestCaseName); @@ -44,6 +45,7 @@ INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_Blocked, ConvertCPULayerTest, ::testing::ValuesIn(inShapes_4D_blocked), ::testing::ValuesIn(precisions()), ::testing::ValuesIn(precisions()), + ::testing::Values(ov::test::SpecialValue::none), ::testing::ValuesIn(filterCPUSpecificParams(memForm4D_static_blocked))), ConvertCPULayerTest::getTestCaseName); @@ -52,6 +54,7 @@ INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_BOOL_Static, ConvertToBoolean ::testing::ValuesIn(inShapes_4D_static()), ::testing::ValuesIn(precisions_floating_point), ::testing::Values(ov::element::boolean), + ::testing::Values(ov::test::SpecialValue::none), ::testing::Values(CPUSpecificParams({nchw}, {nchw}, {}, {}))), ConvertToBooleanCPULayerTest::getTestCaseName); @@ -60,6 +63,7 @@ INSTANTIATE_TEST_SUITE_P(smoke_ConvertCPULayerTest_BOOL_Dynamic, ConvertToBoolea ::testing::ValuesIn(inShapes_4D_dynamic()), ::testing::ValuesIn(precisions_floating_point), ::testing::Values(ov::element::boolean), + ::testing::Values(ov::test::SpecialValue::none), ::testing::Values(CPUSpecificParams({nchw}, {nchw}, {}, {}))), ConvertToBooleanCPULayerTest::getTestCaseName); diff --git a/src/plugins/intel_cpu/tests/functional/shared_tests_instances/single_layer_tests/conversion.cpp b/src/plugins/intel_cpu/tests/functional/shared_tests_instances/single_layer_tests/conversion.cpp index 9ff4d0b989fefa..903b8c083b1a1f 100644 --- a/src/plugins/intel_cpu/tests/functional/shared_tests_instances/single_layer_tests/conversion.cpp +++ b/src/plugins/intel_cpu/tests/functional/shared_tests_instances/single_layer_tests/conversion.cpp @@ -32,6 +32,17 @@ const std::vector types = { ov::element::f64, }; +const std::vector floatTypes = { + ov::element::f32, + ov::element::f16, + ov::element::bf16, +}; + +const std::vector f8Types = { + ov::element::f8e4m3, + ov::element::f8e5m2, +}; + INSTANTIATE_TEST_SUITE_P(smoke_ConversionLayerTest, ConversionLayerTest, ::testing::Combine(::testing::ValuesIn(conversionOpTypes), @@ -49,4 +60,23 @@ INSTANTIATE_TEST_SUITE_P(smoke_ConversionToBooleanLayerTest, ::testing::Values(ov::element::boolean), ::testing::Values(ov::test::utils::DEVICE_CPU)), ConversionLayerTest::getTestCaseName); + +INSTANTIATE_TEST_SUITE_P(smoke_ConversionToF8LayerTest, + ConversionLayerTest, + ::testing::Combine(::testing::Values(conversionOpTypes[0]), + ::testing::ValuesIn(ov::test::static_shapes_to_test_representation(shapes)), + ::testing::ValuesIn(floatTypes), + ::testing::ValuesIn(f8Types), + ::testing::Values(ov::test::utils::DEVICE_CPU)), + ConversionLayerTest::getTestCaseName); + +INSTANTIATE_TEST_SUITE_P(smoke_ConversionFromF8LayerTest, + ConversionLayerTest, + ::testing::Combine(::testing::Values(conversionOpTypes[0]), + ::testing::ValuesIn(ov::test::static_shapes_to_test_representation(shapes)), + ::testing::ValuesIn(f8Types), + ::testing::ValuesIn(floatTypes), + ::testing::Values(ov::test::utils::DEVICE_CPU)), + ConversionLayerTest::getTestCaseName); + } // namespace diff --git a/src/plugins/intel_cpu/tests/functional/shared_tests_instances/single_layer_tests/fake_convert.cpp b/src/plugins/intel_cpu/tests/functional/shared_tests_instances/single_layer_tests/fake_convert.cpp new file mode 100644 index 00000000000000..a2f17ea72cbb3e --- /dev/null +++ b/src/plugins/intel_cpu/tests/functional/shared_tests_instances/single_layer_tests/fake_convert.cpp @@ -0,0 +1,59 @@ +// Copyright (C) 2024 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "single_op_tests/fake_convert.hpp" + +namespace { +using ov::test::FakeConvertLayerTest; + +const std::vector> shapes = {{{2, 3, 4, 5}}}; + +const std::vector data_precisions = {ov::element::f32, ov::element::f16, ov::element::bf16}; + +const std::vector destination_precisions = {ov::element::f8e4m3, ov::element::f8e5m2}; + +const std::vector default_shift = {true, false}; + +const auto simple_fake_convert_params = + ::testing::Combine(::testing::ValuesIn(ov::test::static_shapes_to_test_representation(shapes)), + ::testing::Values(ov::Shape{1}), + ::testing::Values(ov::Shape{1}), + ::testing::ValuesIn(data_precisions), + ::testing::ValuesIn(destination_precisions), + ::testing::ValuesIn(default_shift), + ::testing::Values(ov::test::utils::DEVICE_CPU)); + +const auto broadcast_fake_convert_params = + ::testing::Combine(::testing::ValuesIn(ov::test::static_shapes_to_test_representation(shapes)), + ::testing::Values(ov::Shape{2, 3, 1, 1}), + ::testing::Values(ov::Shape{2, 3, 1, 1}), + ::testing::ValuesIn(data_precisions), + ::testing::ValuesIn(destination_precisions), + ::testing::ValuesIn(default_shift), + ::testing::Values(ov::test::utils::DEVICE_CPU)); + +const auto elementwise_fake_convert_params = + ::testing::Combine(::testing::ValuesIn(ov::test::static_shapes_to_test_representation(shapes)), + ::testing::Values(ov::Shape{2, 3, 4, 5}), + ::testing::Values(ov::Shape{2, 3, 4, 5}), + ::testing::ValuesIn(data_precisions), + ::testing::ValuesIn(destination_precisions), + ::testing::ValuesIn(default_shift), + ::testing::Values(ov::test::utils::DEVICE_CPU)); + +INSTANTIATE_TEST_SUITE_P(smoke_FakeConvert_simple, + FakeConvertLayerTest, + simple_fake_convert_params, + FakeConvertLayerTest::getTestCaseName); + +INSTANTIATE_TEST_SUITE_P(smoke_FakeConvert_broadcast, + FakeConvertLayerTest, + broadcast_fake_convert_params, + FakeConvertLayerTest::getTestCaseName); + +INSTANTIATE_TEST_SUITE_P(smoke_FakeConvert_elementwise, + FakeConvertLayerTest, + elementwise_fake_convert_params, + FakeConvertLayerTest::getTestCaseName); +} // namespace diff --git a/src/plugins/intel_cpu/tests/functional/shared_tests_instances/skip_tests_config.cpp b/src/plugins/intel_cpu/tests/functional/shared_tests_instances/skip_tests_config.cpp index 7af707df602bfc..4c34b3fd2506ac 100644 --- a/src/plugins/intel_cpu/tests/functional/shared_tests_instances/skip_tests_config.cpp +++ b/src/plugins/intel_cpu/tests/functional/shared_tests_instances/skip_tests_config.cpp @@ -173,6 +173,8 @@ std::vector disabledTestPatterns() { R"(.*smoke_TopK/TopKLayerTest.Inference.*_k=21_.*_sort=value_modelType=f16_trgDev=CPU.*)", // Issue: 121812 R"(.*ConvertCPULayerTest.*outFmts=(nhwc|nChw8c|nChw16c).*)", + // Issue: MFDNN-12917. The oneDNN emitter of conversion from fp32 to fp8 has rounding issue. + R"(.*ConvertCPULayerTest.*(\[1.1.1080.1920\]|\(2.17.5.4\))_.*_inputPRC=f32_targetPRC=f8e4m3_.*)", // Need to generate sequence exactly in the i64 data type. Enable in scope of i64 enabling. R"(.*RandomUniformLayerTestCPU.*OutPrc=i64.*)", // Issue: 123815 (Tests are sensintive to available thread count on testing machines) @@ -529,6 +531,7 @@ std::vector disabledTestPatterns() { retVector.emplace_back(R"(.*INFERENCE_PRECISION_HINT=(F|f)16.*)"); retVector.emplace_back(R"(.*ConcatMultiQuerySDPTest.*f16.*)"); retVector.emplace_back(R"(.*ConcatSDPTest.*f16.*)"); + retVector.emplace_back(R"(.*ConvertCPULayerTest.*f16.*)"); } #elif defined(OPENVINO_ARCH_ARM64) || defined(OPENVINO_ARCH_ARM) if (!ov::intel_cpu::hasHardwareSupport(ov::element::f16)) { @@ -536,6 +539,7 @@ std::vector disabledTestPatterns() { retVector.emplace_back(R"(.*INFERENCE_PRECISION_HINT=(F|f)16.*)"); retVector.emplace_back(R"(.*Prc=f16.*)"); retVector.emplace_back(R"(.*ConcatMultiQuerySDPTest.*f16.*HasShapeOf=1.*)"); + retVector.emplace_back(R"(.*ConvertCPULayerTest.*f16.*)"); } else { // Issue 117407 retVector.emplace_back( diff --git a/src/plugins/intel_gpu/include/intel_gpu/primitives/paged_attention.hpp b/src/plugins/intel_gpu/include/intel_gpu/primitives/paged_attention.hpp index f87f608597a6bb..2638f2ad60cf26 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/primitives/paged_attention.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/primitives/paged_attention.hpp @@ -24,6 +24,10 @@ struct paged_attention : public primitive_base { OPENVINO_ASSERT(inputs.size() == 13, "[GPU] Unexpected inputs number for PagedAttention primitive: ", inputs.size()); } + bool has_scores_output() const { + return num_outputs == 2; + } + bool operator==(const primitive& rhs) const override { return compare_common_params(rhs); } diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/paged_attention.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/paged_attention.cpp index 9cf1a252564934..2bc377f2c1459a 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/paged_attention.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/paged_attention.cpp @@ -63,6 +63,7 @@ struct paged_attention_impl : multi_stage_primitive { void load(BinaryInputBuffer& ib) override { parent::load(ib); + ib >> make_data(&has_scores_output, sizeof(bool)); if (is_dynamic()) { auto& kv_cache_update_kernel_selector = kv_cache_update_kernel_selector_t::Instance(); auto kv_cache_update_kernel_impl = kv_cache_update_kernel_selector.GetImplementation(_kernels_data[Stage::KV_CACHE_UPDATE].kernelName); @@ -78,7 +79,45 @@ struct paged_attention_impl : multi_stage_primitive { } } + void save(BinaryOutputBuffer& ob) const override { + parent::save(ob); + ob << make_data(&has_scores_output, sizeof(bool)); + } + std::vector get_internal_buffer_layouts_impl() const override { + /* + * Internal buffers allocation owners and users: + * +--------------------------------------+--------------------+--------------------+ + * | Stage | Allocates & uses | Reuses | + * +--------------------------------------+--------------------+--------------------+ + * | KV_CACHE_UPDATE | [0, 1, 2] | | + * +--------------------------------------+--------------------+--------------------+ + * | SDPA (1st token) | | [0, 1, 2] | + * +--------------------------------------+--------------------+--------------------+ + * | PA_SDPA (2nd+ token) | [5, 6, 7] | | + * +--------------------------------------+--------------------+--------------------+ + * | PA_SDPA (mixed mode) | [5, 6, 7, 8] | | + * +--------------------------------------+--------------------+--------------------+ + * | SDPA (1st token) + scores output | | [0, 1, 2, 3, 4] | + * +--------------------------------------+--------------------+--------------------+ + * | PA_SDPA (2nd+ token) + scores output | [3, 4, 5, 6, 7] | | + * +--------------------------------------+--------------------+--------------------+ + * | PA_SDPA (mixed mode) + scores output | [3, 4, 5, 6, 7, 8] | | + * +--------------------------------------+--------------------+--------------------+ + * + * Description: + * 0, 1, 2 - Buffers used for proper blocks distribution for kv_cache_update and + * sdpa_opt (1st token calculation) block configuration over target_seq_len dimension. + * Filled in paged_attention_inst::on_execute() call. + * 3, 4 - Optional buffers used for PA scores output calculation, storing intermediate + * softmax values by partitions (filled in PA/SDPA kernels) and sequence length offsets + * for each subsequence (filled in paged_attention_inst::on_execute() call). + * 5, 6, 7 - Used for 2nd+ PA calculation (for softmax exp_sums, max_logits, and intermediate output). + * Filled in PA/SDPA kernels. + * 8 - Optional buffer used for mixed PA execution mode, mapping gws idx to subsequence id. + * Filled in paged_attention_inst::on_execute() call. + */ + auto add_internal_buffers = [](std::vector& layouts, const kernel_selector::KernelData& kd) { if (kd.internalBufferSizes.empty()) return; @@ -133,6 +172,7 @@ struct paged_attention_impl : multi_stage_primitive { args.outputs = { instance.output_memory_ptr(0) }; } else if (stage == Stage::PA_SDPA) { if (kernel_idx == 0 || kernel_idx == 1) { + // 2nd+ token calculation or mixed stage tokens calculation args.shape_info = instance.shape_info_memory_ptr(); args.inputs = { instance.input_memory_ptr(0), @@ -155,7 +195,8 @@ struct paged_attention_impl : multi_stage_primitive { if (desc->has_alibi) { args.inputs.push_back(instance.alibi_memory_ptr()); } - } else { + } else if (kernel_idx == 2 || kernel_idx == 3) { + // Finalization kernel or mixed stage finalization kernel args.inputs = { instance.past_lens_memory_ptr() }; if (is_mixed_mode) { @@ -163,17 +204,31 @@ struct paged_attention_impl : multi_stage_primitive { // dependency args.inputs.push_back(instance.subsequence_begins_memory_ptr()); } + } else if (kernel_idx == 4) { + // Output scores calculation kernel + args.inputs = { instance.past_lens_memory_ptr(), + instance.subsequence_begins_memory_ptr() }; } args.outputs = { instance.output_memory_ptr(0) }; + + if (kernel_idx == 4) { + args.outputs.push_back(instance.output_memory_ptr(1)); + } } return args; } std::set get_lockable_internal_buffers() const override { - return std::set{ 0, 1, 2, /* SDPA and KV_CACHE_UPDATE indexes configuration */ - 6, /* PA_SDPA multiple tokens mode */ }; + size_t mixed_mode_buffer = has_scores_output ? 8 : 6; + + std::set lockable_ids = { 0, 1, 2, /* SDPA and KV_CACHE_UPDATE indexes configuration */ + mixed_mode_buffer /* PA_SDPA multiple tokens mode */ }; + if (has_scores_output) + lockable_ids.insert(4 /* Precalculated accumulated sequence length offsets for each subsequence */); + + return lockable_ids; }; void execute_stage(const std::vector& events, @@ -194,8 +249,17 @@ struct paged_attention_impl : multi_stage_primitive { if (stage == Stage::PA_SDPA) { internal_buffers_offset = _kernels_data[Stage::KV_CACHE_UPDATE].internalBufferSizes.size(); internal_buffers_count = _kernels_data[Stage::PA_SDPA].internalBufferSizes.size(); - } else { + } else if (stage == Stage::KV_CACHE_UPDATE) { + internal_buffers_count = _kernels_data[Stage::KV_CACHE_UPDATE].internalBufferSizes.size(); + } else if (stage == Stage::SDPA) { internal_buffers_count = _kernels_data[Stage::KV_CACHE_UPDATE].internalBufferSizes.size(); + + const auto desc = instance.get_node().as().get_primitive(); + if (desc->has_scores_output()) { + // Add intermediate buffers for PagedAttention scores calculation: + // softmax_results, subsequence_offsets, exp_sums, max_logits, tmp_out + internal_buffers_count += 5; + } } for (size_t kd_idx = 0; kd_idx < _kernels_data[stage].kernels.size(); ++kd_idx) { @@ -216,6 +280,23 @@ struct paged_attention_impl : multi_stage_primitive { intermediate_memories.begin() + internal_buffers_offset, intermediate_memories.begin() + internal_buffers_offset + internal_buffers_count); + GPU_DEBUG_TRACE_DETAIL << "Execute stage=" << stage << " kernel=" << kd_idx << " " << _kernels_data[stage].kernelName << " start_offset=" + << internal_buffers_offset << " count=" << internal_buffers_count << "\n"; + + GPU_DEBUG_TRACE_DETAIL << "Configured kernel arguments:\n"; + for (size_t i = 0; i < _kernels_data[stage].kernels[kd_idx].params.arguments.size(); i++) { + GPU_DEBUG_TRACE_DETAIL << "\t" << i << ": type=" << static_cast(_kernels_data[stage].kernels[kd_idx].params.arguments[i].t) << " " + << "index=" << _kernels_data[stage].kernels[kd_idx].params.arguments[i].index << "\n"; + } + + GPU_DEBUG_TRACE_DETAIL << "Memory buffers:" + << "shape_info=" << args.shape_info << " " + << "inputs=" << args.inputs.size() << " " + << "outputs=" << args.outputs.size() << " " + << "intermediates=" << args.intermediates.size() << " " + << "weights=" << args.weights << " " + << "scalars=" << (args.scalars ? args.scalars->size() : 0) << "\n"; + stream.set_arguments(*_kernels[idx_final], _kernels_data[stage].kernels[kd_idx].params, args); const auto& gws = params.workGroups.global; @@ -242,10 +323,13 @@ struct paged_attention_impl : multi_stage_primitive { execute_stage(events, instance, res_events, Stage::KV_CACHE_UPDATE, is_mixed_mode); - std::vector dep_events(res_events.begin(), res_events.end()); if (stage == PagedAttentionStage::PREFILL) { + std::vector dep_events(res_events.begin(), res_events.end()); execute_stage(dep_events, instance, res_events, Stage::SDPA, is_mixed_mode); - } else if (stage == PagedAttentionStage::GENERATE || stage == PagedAttentionStage::MIXED) { + } + + if (stage == PagedAttentionStage::GENERATE || stage == PagedAttentionStage::MIXED || has_scores_output) { + std::vector dep_events(res_events.begin(), res_events.end()); execute_stage(dep_events, instance, res_events, Stage::PA_SDPA, is_mixed_mode); } @@ -338,7 +422,7 @@ struct paged_attention_impl : multi_stage_primitive { return aligned_seq_len; } - static kernel_selector::sdpa_configuration get_sdpa_configuration(const kernel_impl_params& impl_param) { + static kernel_selector::sdpa_configuration get_sdpa_configuration(const kernel_impl_params& impl_param, bool is_dynamic = true) { kernel_selector::sdpa_configuration config; const auto desc = impl_param.typed_desc(); @@ -362,37 +446,45 @@ struct paged_attention_impl : multi_stage_primitive { config.group_size = desc->heads_num / desc->kv_heads_num; } + if (desc->has_scores_output() && !is_dynamic) { + const auto& input_mem = impl_param.memory_deps; + const auto max_context_len = input_mem.at(12); + mem_lock max_context_len_mem_lock(max_context_len, *impl_param.strm); + config.paged_attention_max_len = max_context_len_mem_lock[0]; + } + return config; } static kv_cache_update_kernel_params_t get_kv_cache_update_kernel_params(const kernel_impl_params& impl_param, const PagedAttentionStage& stage, + const kernel_selector::MultiDataTensor& input_tensors, bool is_dynamic = false) { auto params = get_default_params(impl_param, is_dynamic); - const auto& key_layout = impl_param.get_input_layout(1); - const auto& value_layout = impl_param.get_input_layout(2); - const auto& key_cache_layout = impl_param.get_input_layout(3); - const auto& value_cache_layout = impl_param.get_input_layout(4); - const auto& past_lens_layout = impl_param.get_input_layout(5); - const auto& block_indices_layout = impl_param.get_input_layout(7); - const auto& block_indices_begins_layout = impl_param.get_input_layout(8); - const auto& subsequence_begins_layout = impl_param.get_input_layout(6); + const auto& key_tensor = input_tensors[1]; + const auto& value_tensor = input_tensors[2]; + const auto& key_cache_tensor = input_tensors[3]; + const auto& value_cache_tensor = input_tensors[4]; + const auto& past_lens_tensor = input_tensors[5]; + const auto& block_indices_tensor = input_tensors[7]; + const auto& block_indices_begins_tensor = input_tensors[8]; + const auto& subsequence_begins_tensor = input_tensors[6]; const auto inputs_number = 6; const auto outputs_number = 2; params.inputs.resize(inputs_number); params.outputs.resize(outputs_number); - params.inputs[0] = convert_data_tensor(key_layout); - params.inputs[1] = convert_data_tensor(value_layout); - params.inputs[2] = convert_data_tensor(past_lens_layout); - params.inputs[3] = convert_data_tensor(block_indices_layout); - params.inputs[4] = convert_data_tensor(block_indices_begins_layout); - params.inputs[5] = convert_data_tensor(subsequence_begins_layout); - params.outputs[0] = convert_data_tensor(key_cache_layout); - params.outputs[1] = convert_data_tensor(value_cache_layout); + params.inputs[0] = key_tensor; + params.inputs[1] = value_tensor; + params.inputs[2] = past_lens_tensor; + params.inputs[3] = block_indices_tensor; + params.inputs[4] = block_indices_begins_tensor; + params.inputs[5] = subsequence_begins_tensor; + params.outputs[0] = key_cache_tensor; + params.outputs[1] = value_cache_tensor; - params.conf = get_sdpa_configuration(impl_param); + params.conf = get_sdpa_configuration(impl_param, is_dynamic); params.is_prefill = stage == PagedAttentionStage::PREFILL || stage == PagedAttentionStage::MIXED; @@ -418,18 +510,23 @@ struct paged_attention_impl : multi_stage_primitive { return params; } - static sdpa_kernel_params_t get_sdpa_kernel_params(const kernel_impl_params& impl_param, const PagedAttentionStage& stage, bool is_dynamic = false) { + static sdpa_kernel_params_t get_sdpa_kernel_params(const kernel_impl_params& impl_param, + const PagedAttentionStage& stage, + const kernel_selector::MultiDataTensor& input_tensors, + bool is_dynamic = false) { const auto desc = impl_param.typed_desc(); auto params = get_default_params(impl_param, is_dynamic); - const auto& query_layout = impl_param.get_input_layout(0); - const auto& key_layout = impl_param.get_input_layout(1); - const auto& value_layout = impl_param.get_input_layout(2); - const auto& subsequence_begins_layout = impl_param.get_input_layout(6); - const auto& scale_layout = impl_param.get_input_layout(9); - const auto& alibi_layout = impl_param.get_input_layout(11); - const auto has_alibi = alibi_layout.count() > 0; + const auto& query_tensor = input_tensors[0]; + const auto& key_tensor = input_tensors[1]; + const auto& value_tensor = input_tensors[2]; + const auto& subsequence_begins_tensor = input_tensors[6]; + const auto& scale_tensor = input_tensors[9]; + const auto& alibi_tensor = input_tensors[11]; + + const auto has_alibi = impl_param.get_input_layout(11).count() > 0; const auto has_scale_input = !desc->scale_val.has_value(); + const auto has_scores_output = desc->has_scores_output(); auto inputs_number = 4; if (has_scale_input) @@ -440,18 +537,23 @@ struct paged_attention_impl : multi_stage_primitive { auto input_idx = 0; params.inputs.resize(inputs_number); - params.inputs[input_idx++] = convert_data_tensor(query_layout); - params.inputs[input_idx++] = convert_data_tensor(key_layout); - params.inputs[input_idx++] = convert_data_tensor(value_layout); - params.inputs[input_idx++] = convert_data_tensor(subsequence_begins_layout); + params.inputs[input_idx++] = query_tensor; + params.inputs[input_idx++] = key_tensor; + params.inputs[input_idx++] = value_tensor; + params.inputs[input_idx++] = subsequence_begins_tensor; if (has_scale_input) - params.inputs[input_idx++] = convert_data_tensor(scale_layout); + params.inputs[input_idx++] = scale_tensor; if (has_alibi) - params.inputs[input_idx++] = convert_data_tensor(alibi_layout); + params.inputs[input_idx++] = alibi_tensor; - params.conf = get_sdpa_configuration(impl_param); + if (has_scores_output) { + params.outputs.resize(2); + params.outputs[1] = convert_data_tensor(impl_param.get_output_layout(1)); + } + + params.conf = get_sdpa_configuration(impl_param, is_dynamic); const auto& in_offsets_map = impl_param.in_port_to_shape_info_offset; const auto& out_offsets_map = impl_param.out_port_to_shape_info_offset; @@ -475,26 +577,34 @@ struct paged_attention_impl : multi_stage_primitive { if ((stage == PagedAttentionStage::PREFILL || stage == PagedAttentionStage::MIXED) && !is_dynamic) params.conf.paged_attention_aligned_seq_len = get_aligned_seq_len(impl_param, stage); + if (has_scores_output) + out_tensor_to_offset_map.insert({1, out_offsets_map.at(1)}); + params.set_dynamic_shape_offsets(in_tensor_to_offset_map, out_tensor_to_offset_map); return params; } - static pa_sdpa_kernel_params_t get_pa_sdpa_params(const kernel_impl_params& impl_param, const PagedAttentionStage& stage, bool is_dynamic = false) { + static pa_sdpa_kernel_params_t get_pa_sdpa_params(const kernel_impl_params& impl_param, + const PagedAttentionStage& stage, + const kernel_selector::MultiDataTensor& input_tensors, + bool is_dynamic = false) { const auto desc = impl_param.typed_desc(); auto params = get_default_params(impl_param, is_dynamic); - const auto& query_layout = impl_param.get_input_layout(0); - const auto& key_cache_layout = impl_param.get_input_layout(3); - const auto& value_cache_layout = impl_param.get_input_layout(4); - const auto& past_lens_layout = impl_param.get_input_layout(5); - const auto& block_indices_layout = impl_param.get_input_layout(7); - const auto& block_indices_begins_layout = impl_param.get_input_layout(8); - const auto& subsequence_begins_layout = impl_param.get_input_layout(6); - const auto& scale_layout = impl_param.get_input_layout(9); - const auto& alibi_layout = impl_param.get_input_layout(11); - const auto has_alibi = alibi_layout.count() > 0; + const auto& query_tensor = input_tensors[0]; + const auto& key_cache_tensor = input_tensors[3]; + const auto& value_cache_tensor = input_tensors[4]; + const auto& past_lens_tensor = input_tensors[5]; + const auto& block_indices_tensor = input_tensors[7]; + const auto& block_indices_begins_tensor = input_tensors[8]; + const auto& subsequence_begins_tensor = input_tensors[6]; + const auto& scale_tensor = input_tensors[9]; + const auto& alibi_tensor = input_tensors[11]; + + const auto has_alibi = impl_param.get_input_layout(11).count() > 0; const auto has_scale_input = !desc->scale_val.has_value(); + const auto has_scores_output = desc->has_scores_output(); auto inputs_number = 7; if (has_scale_input) @@ -505,28 +615,34 @@ struct paged_attention_impl : multi_stage_primitive { auto input_idx = 0; params.inputs.resize(inputs_number); - params.inputs[input_idx++] = convert_data_tensor(query_layout); - params.inputs[input_idx++] = convert_data_tensor(key_cache_layout); - params.inputs[input_idx++] = convert_data_tensor(value_cache_layout); - params.inputs[input_idx++] = convert_data_tensor(past_lens_layout); - params.inputs[input_idx++] = convert_data_tensor(block_indices_layout); - params.inputs[input_idx++] = convert_data_tensor(block_indices_begins_layout); - params.inputs[input_idx++] = convert_data_tensor(subsequence_begins_layout); - params.conf = get_sdpa_configuration(impl_param); + params.inputs[input_idx++] = query_tensor; + params.inputs[input_idx++] = key_cache_tensor; + params.inputs[input_idx++] = value_cache_tensor; + params.inputs[input_idx++] = past_lens_tensor; + params.inputs[input_idx++] = block_indices_tensor; + params.inputs[input_idx++] = block_indices_begins_tensor; + params.inputs[input_idx++] = subsequence_begins_tensor; + + params.conf = get_sdpa_configuration(impl_param, is_dynamic); if (has_scale_input) - params.inputs[input_idx++] = convert_data_tensor(scale_layout); + params.inputs[input_idx++] = scale_tensor; if (has_alibi) - params.inputs[input_idx++] = convert_data_tensor(alibi_layout); + params.inputs[input_idx++] = alibi_tensor; - params.multi_tokens_mode = stage == PagedAttentionStage::MIXED; + if (has_scores_output) { + params.outputs.resize(2); + params.outputs[1] = convert_data_tensor(impl_param.get_output_layout(1)); + } - if ((stage == PagedAttentionStage::GENERATE || stage == PagedAttentionStage::MIXED) && !is_dynamic) { + params.stage = stage; + + if (!has_scores_output && !is_dynamic) { const auto& input_mem = impl_param.memory_deps; const auto max_context_len = input_mem.at(12); mem_lock max_context_len_mem_lock(max_context_len, *impl_param.strm); - params.max_context_len = max_context_len_mem_lock[0]; + params.conf.paged_attention_max_len = max_context_len_mem_lock[0]; } const auto& in_offsets_map = impl_param.in_port_to_shape_info_offset; @@ -552,6 +668,9 @@ struct paged_attention_impl : multi_stage_primitive { if (has_alibi) in_tensor_to_offset_map.insert({input_idx++, in_offsets_map.at(11)}); + if (has_scores_output) + out_tensor_to_offset_map.insert({1, out_offsets_map.at(1)}); + params.set_dynamic_shape_offsets(in_tensor_to_offset_map, out_tensor_to_offset_map); return params; @@ -560,14 +679,20 @@ struct paged_attention_impl : multi_stage_primitive { void update_dispatch_data(const kernel_impl_params& impl_param) override { const auto stage = get_paged_attention_stage(impl_param); - auto kv_cache_update_kernel_params = get_kv_cache_update_kernel_params(impl_param, stage, impl_param.is_dynamic()); + kernel_selector::MultiDataTensor input_tensors; + for (const auto& input_layout : impl_param.input_layouts) + input_tensors.emplace_back(convert_data_tensor(input_layout)); + + auto kv_cache_update_kernel_params = get_kv_cache_update_kernel_params(impl_param, stage, input_tensors, impl_param.is_dynamic()); (_kernels_data[Stage::KV_CACHE_UPDATE].update_dispatch_data_func)(kv_cache_update_kernel_params, _kernels_data[Stage::KV_CACHE_UPDATE]); if (stage == PagedAttentionStage::PREFILL) { - auto sdpa_kernel_params = get_sdpa_kernel_params(impl_param, stage, impl_param.is_dynamic()); + auto sdpa_kernel_params = get_sdpa_kernel_params(impl_param, stage, input_tensors, impl_param.is_dynamic()); (_kernels_data[Stage::SDPA].update_dispatch_data_func)(sdpa_kernel_params, _kernels_data[Stage::SDPA]); - } else if (stage == PagedAttentionStage::GENERATE || stage == PagedAttentionStage::MIXED) { - auto pa_sdpa_kernel_params = get_pa_sdpa_params(impl_param, stage, impl_param.is_dynamic()); + } + + if (stage == PagedAttentionStage::GENERATE || stage == PagedAttentionStage::MIXED || has_scores_output) { + auto pa_sdpa_kernel_params = get_pa_sdpa_params(impl_param, stage, input_tensors, impl_param.is_dynamic()); (_kernels_data[Stage::PA_SDPA].update_dispatch_data_func)(pa_sdpa_kernel_params, _kernels_data[Stage::PA_SDPA]); } } @@ -576,20 +701,32 @@ struct paged_attention_impl : multi_stage_primitive { std::vector kernels_data; const auto stage = PagedAttentionStage::UNKNOWN; - auto kv_cache_update_kernel_params = get_kv_cache_update_kernel_params(impl_param, stage, impl_param.is_dynamic()); + kernel_selector::MultiDataTensor input_tensors; + for (const auto& input_layout : impl_param.input_layouts) + input_tensors.emplace_back(convert_data_tensor(input_layout)); + + auto kv_cache_update_kernel_params = get_kv_cache_update_kernel_params(impl_param, stage, input_tensors, impl_param.is_dynamic()); auto& kv_cache_update_kernel_selector = kv_cache_update_kernel_selector_t::Instance(); kernels_data.push_back(kv_cache_update_kernel_selector.get_best_kernel(kv_cache_update_kernel_params)); - auto sdpa_kernel_params = get_sdpa_kernel_params(impl_param, stage, impl_param.is_dynamic()); + auto sdpa_kernel_params = get_sdpa_kernel_params(impl_param, stage, input_tensors, impl_param.is_dynamic()); auto& sdpa_kernel_selector = sdpa_kernel_selector_t::Instance(); kernels_data.push_back(sdpa_kernel_selector.get_best_kernel(sdpa_kernel_params)); - auto pa_sdpa_kernel_params = get_pa_sdpa_params(impl_param, stage, impl_param.is_dynamic()); + auto pa_sdpa_kernel_params = get_pa_sdpa_params(impl_param, stage, input_tensors, impl_param.is_dynamic()); auto& pa_sdpa_kernel_selector = pa_sdpa_kernel_selector_t::Instance(); kernels_data.push_back(pa_sdpa_kernel_selector.get_best_kernel(pa_sdpa_kernel_params)); - return cldnn::make_unique(kernels_data); + auto impl = cldnn::make_unique(kernels_data); + + const auto& desc = impl_param.typed_desc(); + impl->has_scores_output = desc->has_scores_output(); + + return impl; } + +private: + bool has_scores_output = false; }; namespace detail { diff --git a/src/plugins/intel_gpu/src/graph/include/paged_attention_inst.h b/src/plugins/intel_gpu/src/graph/include/paged_attention_inst.h index a7918ba9c3719c..675d77296aa06b 100644 --- a/src/plugins/intel_gpu/src/graph/include/paged_attention_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/paged_attention_inst.h @@ -7,14 +7,11 @@ #include "intel_gpu/primitives/paged_attention.hpp" #include "primitive_inst.h" +#include "sdpa/pa_sdpa_kernel_opt.h" + namespace cldnn { -enum PagedAttentionStage { - GENERATE = 0, - PREFILL = 1, - MIXED = 2, - UNKNOWN = 3 -}; +using PagedAttentionStage = kernel_selector::PagedAttentionStage; PagedAttentionStage get_paged_attention_stage(const kernel_impl_params& impl_param); @@ -61,6 +58,9 @@ class typed_primitive_inst : public typed_primitive_inst_base

prefill_network; diff --git a/src/plugins/intel_gpu/src/graph/paged_attention.cpp b/src/plugins/intel_gpu/src/graph/paged_attention.cpp index 787fd184f75b6a..c761aaf63799cd 100644 --- a/src/plugins/intel_gpu/src/graph/paged_attention.cpp +++ b/src/plugins/intel_gpu/src/graph/paged_attention.cpp @@ -48,14 +48,38 @@ layout paged_attention_inst::calc_output_layout(const paged_attention_node& /*no template std::vector paged_attention_inst::calc_output_layouts(paged_attention_node const& /*node*/, kernel_impl_params const& impl_param) { - auto out_layout = impl_param.get_input_layout(0); + auto data_layout = impl_param.get_input_layout(0); const auto& key_cache_ps = impl_param.get_input_layout(3).get_partial_shape(); bool valid_block_size = key_cache_ps[3].is_dynamic() || key_cache_ps[3].get_length() == paged_attention::block_size; OPENVINO_ASSERT(valid_block_size, "[GPU] Incorrect block size for Paged Attention operation. " "Expected ", paged_attention::block_size, ", but got ", key_cache_ps[3].get_length()); - return {out_layout}; + std::vector output_layouts{ data_layout }; + + const auto& desc = impl_param.typed_desc(); + if (desc->has_scores_output()) { + const auto past_lens_idx = 5; + const auto output_dt = data_layout.data_type; + if (impl_param.get_input_layout(past_lens_idx).is_static()) { + const auto& memory_deps = impl_param.memory_deps; + const auto past_lens_mem = memory_deps.at(past_lens_idx); + mem_lock past_lens_mem_lock(past_lens_mem, *impl_param.strm); + + long int total_size = 0; + for (size_t i = 0; i < past_lens_mem_lock.size(); i++) { + total_size += past_lens_mem_lock[i]; + } + + total_size += static_cast(impl_param.get_input_layout(0).get_shape()[0]); + + output_layouts.push_back(layout{ov::PartialShape{total_size}, output_dt, format::bfyx}); + } else { + output_layouts.push_back(layout{ov::PartialShape::dynamic(1), output_dt, format::bfyx}); + } + } + + return output_layouts; } template std::vector @@ -81,45 +105,79 @@ std::string paged_attention_inst::to_string(const paged_attention_node& node) { } void paged_attention_inst::on_execute() { - auto stage = get_paged_attention_stage(*_impl_params); + const auto& desc = _impl_params->typed_desc(); + const bool has_scores_output = desc->has_scores_output(); + const auto stage = get_paged_attention_stage(*_impl_params); - if (stage == PagedAttentionStage::UNKNOWN || - stage == PagedAttentionStage::GENERATE) + if ((stage == PagedAttentionStage::UNKNOWN) || + (stage == PagedAttentionStage::GENERATE && !has_scores_output)) return; + auto& stream = get_network().get_stream(); + const auto past_lens_mem = past_lens_memory_ptr(); + const auto subsequence_begins_mem = subsequence_begins_memory_ptr(); + mem_lock past_lens_mem_lock(past_lens_mem, stream); + mem_lock subsequence_begins_mem_lock(subsequence_begins_mem, stream); + std::unique_ptr> subsequence_offsets_lock = nullptr; + + if (has_scores_output) { + const size_t subsequence_offsets_idx = 4; + + OPENVINO_ASSERT(_intermediates_memory.size() > subsequence_offsets_idx, + "[GPU] Unexpected number of intermediates buffers for Paged Attention for scores output calculation"); + + auto subsequence_offsets_mem = _intermediates_memory[subsequence_offsets_idx]; + subsequence_offsets_lock.reset(new mem_lock(subsequence_offsets_mem, stream)); + } + + if (stage == PagedAttentionStage::GENERATE) { + // For the generate stage it's not necessary to configure any other intermediate + // buffers. Simply calculate the offsets and exit + size_t subsequence_offsets_acc = 0; + for (size_t i = 0; i < subsequence_begins_mem_lock.size() - 1; i++) { + const auto past_len = past_lens_mem_lock[i]; + const auto seq_start = subsequence_begins_mem_lock[i]; + const auto seq_end = subsequence_begins_mem_lock[i + 1]; + const auto seq_length = seq_end - seq_start; + + if (subsequence_offsets_lock) { + subsequence_offsets_lock->operator[](i) = static_cast(subsequence_offsets_acc); + subsequence_offsets_acc += seq_length + past_len; + } + } + + return; + } + OPENVINO_ASSERT(_intermediates_memory.size() >= 3, "Unexpected number of intermediates buffers for Paged Attention at prefill stage"); const auto blocks_indexes_start_idx = 0; const auto blocks_indexes_end_idx = 1; const auto blocked_gws_subseq_mapping_idx = 2; - const auto past_lens_mem = past_lens_memory_ptr(); - auto subsequence_begins_mem = subsequence_begins_memory_ptr(); auto blocks_indexes_start_mem = _intermediates_memory[blocks_indexes_start_idx]; auto blocks_indexes_end_mem = _intermediates_memory[blocks_indexes_end_idx]; auto blocked_gws_subseq_mapping_mem = _intermediates_memory[blocked_gws_subseq_mapping_idx]; OPENVINO_ASSERT(subsequence_begins_mem->get_layout().data_type == data_types::i32); - auto& stream = get_network().get_stream(); - mem_lock past_lens_mem_lock(past_lens_mem, stream); - mem_lock subsequence_begins_mem_lock(subsequence_begins_mem, stream); mem_lock blocks_indexes_start_lock(blocks_indexes_start_mem, stream); mem_lock blocks_indexes_end_lock(blocks_indexes_end_mem, stream); mem_lock blocked_gws_subseq_mapping_mem_lock(blocked_gws_subseq_mapping_mem, stream); std::unique_ptr> sequential_gws_subseq_mapping_lock = nullptr; if (stage == PagedAttentionStage::MIXED) { - const auto sequential_gws_subseq_mapping_idx = 6; + const size_t sequential_gws_subseq_mapping_idx = has_scores_output ? 8 : 6; OPENVINO_ASSERT(_intermediates_memory.size() > sequential_gws_subseq_mapping_idx, - "Unexpected number of intermediates buffers for Paged Attention for mixed stage"); + "[GPU] Unexpected number of intermediates buffers for Paged Attention for mixed stage"); auto sequential_gws_subseq_mapping_mem = _intermediates_memory[sequential_gws_subseq_mapping_idx]; sequential_gws_subseq_mapping_lock.reset(new mem_lock(sequential_gws_subseq_mapping_mem, stream)); } size_t index = 0; + size_t subsequence_offsets_acc = 0; const auto target_seq_len_block_size = 16; // TODO: Get block size from the impl for (size_t i = 0; i < subsequence_begins_mem_lock.size() - 1; i++) { const auto past_len = past_lens_mem_lock[i]; @@ -159,6 +217,11 @@ void paged_attention_inst::on_execute() { sequential_gws_subseq_mapping_lock->operator[](idx) = static_cast(i); } } + + if (subsequence_offsets_lock) { + subsequence_offsets_lock->operator[](i) = static_cast(subsequence_offsets_acc); + subsequence_offsets_acc += seq_length + past_len; + } } } diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/pa_sdpa_opt.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/pa_sdpa_opt.cl index 00c43829d02ea7..7e960afa4b87d3 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/pa_sdpa_opt.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/pa_sdpa_opt.cl @@ -44,6 +44,10 @@ KERNEL(pa_sdpa_opt)( const __global ALIBI_INPUT_TYPE* alibi_slopes, #endif __global OUTPUT_TYPE* output, +#if PAGED_ATTENTION_SCORES_OUTPUT + __global SOFTMAX_ACCUMULATOR_TYPE* softmax_results, + const __global int* subsequence_offsets, +#endif __global SOFTMAX_ACCUMULATOR_TYPE* exp_sums, __global SOFTMAX_ACCUMULATOR_TYPE* max_logits, __global OUTPUT_TYPE* tmp_out @@ -276,6 +280,28 @@ KERNEL(pa_sdpa_opt)( const uint max_logits_offset = exp_sums_offset; max_logits[max_logits_offset] = qk_max; } + +#if PAGED_ATTENTION_SCORES_OUTPUT +#if MULTI_TOKENS_PROCESSING + const uint subsequence_idx = gws_subseq_mapping[seq_idx]; + const uint subsequence_start_pos = subsequence_begins[subsequence_idx]; + const uint subsequence_end_pos = subsequence_begins[subsequence_idx + 1]; + const bool save_softmax_results = seq_idx == subsequence_end_pos - 1; +#else + const uint subsequence_idx = seq_idx; + const bool save_softmax_results = true; +#endif // MULTI_TOKENS_PROCESSING + // PagedAttention is supposed to save only last "row" of the QK matrix multiplication, + // so save SEQ_LEN_PARTITION_SIZE elements for each partition + if (save_softmax_results) { + const uint output_offset = subsequence_idx * HEADS_NUM * total_partitions_num * SEQ_LEN_PARTITION_SIZE + + head_num_idx * total_partitions_num * SEQ_LEN_PARTITION_SIZE + + partition_idx * SEQ_LEN_PARTITION_SIZE; + for (uint i = sgid * SUBGROUP_SIZE + sglid; i < SEQ_LEN_PARTITION_SIZE; i += SUBGROUPS_PER_WG * SUBGROUP_SIZE) { + softmax_results[output_offset + i] = slm_qk_vals[i]; + } + } +#endif // PAGED_ATTENTION_SCORES_OUTPUT } } @@ -370,6 +396,10 @@ KERNEL(pa_sdpa_finalization_stage)( const __global INPUT6_TYPE* subsequence_begins, #endif __global OUTPUT_TYPE* output, +#if PAGED_ATTENTION_SCORES_OUTPUT + __global SOFTMAX_ACCUMULATOR_TYPE* softmax_results, + const __global int* subsequence_offsets, +#endif const __global SOFTMAX_ACCUMULATOR_TYPE* exp_sums, const __global SOFTMAX_ACCUMULATOR_TYPE* max_logits, const __global OUTPUT_TYPE* tmp_out, @@ -500,3 +530,155 @@ KERNEL(pa_sdpa_finalization_stage)( } #endif + +#ifdef SDPA_STAGE_2 +#define MAX_PARTITIONS_NUM 128 + +REQD_SUB_GROUP_SIZE(SUBGROUP_SIZE) +KERNEL(pa_sdpa_scores_calculation)( + const __global INPUT3_TYPE* past_lens, + const __global INPUT6_TYPE* subsequence_begins, + __global OUTPUT1_TYPE* scores_output, + const __global SOFTMAX_ACCUMULATOR_TYPE* softmax_output, + const __global int* subsequence_offsets, + const __global SOFTMAX_ACCUMULATOR_TYPE* exp_sums, + const __global SOFTMAX_ACCUMULATOR_TYPE* max_logits, + const __global OUTPUT_TYPE* tmp_out, + const uint is_mixed_mode) { + const uint subsequence_idx = get_global_id(2); + const uint partition_global_idx = get_global_id(0); + const uint local_id = get_local_id(0); + const uint partition_idx = get_group_id(0); + const uint partition_size = get_local_size(0); + const uint max_seq_len = get_global_size(0); + const uint partitions_num = get_num_groups(0); + const uint sgid = get_sub_group_id(); + const uint sgid_num = get_num_sub_groups(); + const uint sglid = get_sub_group_local_id(); + + const int subsequence_begin = subsequence_begins[subsequence_idx]; + const int subsequence_end = subsequence_begins[subsequence_idx + 1]; + const uint seq_len = (subsequence_end - subsequence_begin) + past_lens[subsequence_idx]; + + const uint num_of_partitions = CEIL_DIV(seq_len, partition_size); + + if (partition_idx >= num_of_partitions) + return; + + __local SOFTMAX_ACCUMULATOR_TYPE slm_exp_sums[HEADS_NUM]; + __local SOFTMAX_ACCUMULATOR_TYPE slm_global_exp_sum[HEADS_NUM]; + + SOFTMAX_ACCUMULATOR_TYPE total_score = SOFTMAX_ACCUMULATOR_VAL_ZERO; + if (seq_len <= partition_size) { + // If seq_len is less than the partition size, just reduce the results over the heads + for (uint head_idx = 0; head_idx < HEADS_NUM; head_idx++) { + const uint input_offset = subsequence_idx * HEADS_NUM * max_seq_len + head_idx * max_seq_len + partition_global_idx; + SOFTMAX_ACCUMULATOR_TYPE softmax_value = softmax_output[input_offset]; + total_score += softmax_value; + } + } else if (seq_len <= partition_size * MAX_PARTITIONS_NUM) { + // Optimized version for longer prompts (up to partition_size * MAX_PARTITIONS_NUM, ~64K tokens) + + // Depending on the previous kernel exp_sums and max_logits might have different structure: + // For ordinary 1st and 2nd token kernels, there is only a single entry per subsequence. + // However, for mixed mode execution, exp_sums and max_logits include information for all + // tokens of each subsequence, but only the last one is needed for score calculation. + const uint subsequence_pos = is_mixed_mode ? subsequence_end - 1 : subsequence_idx; + + for (uint head_idx = sgid; head_idx < HEADS_NUM; head_idx += sgid_num) { + SOFTMAX_ACCUMULATOR_TYPE max_logit[MAX_PARTITIONS_NUM / SUBGROUP_SIZE]; + SOFTMAX_ACCUMULATOR_TYPE exp_sum[MAX_PARTITIONS_NUM / SUBGROUP_SIZE]; + + const uint exp_sums_offset = subsequence_pos * HEADS_NUM * partitions_num + head_idx * partitions_num; + for (int i = 0; i < partitions_num / SUBGROUP_SIZE; i++) { + max_logit[i] = max_logits[exp_sums_offset + i * SUBGROUP_SIZE + sglid]; + exp_sum[i] = exp_sums[exp_sums_offset + i * SUBGROUP_SIZE + sglid]; + } + + const uint partitions_leftovers = partitions_num % SUBGROUP_SIZE; + if (partitions_leftovers != 0) { + const uint idx = partitions_num / SUBGROUP_SIZE; + max_logit[idx] = sglid >= partitions_leftovers ? SOFTMAX_ACCUMULATOR_VAL_MIN : max_logits[exp_sums_offset + idx * SUBGROUP_SIZE + sglid]; + exp_sum[idx] = sglid >= partitions_leftovers ? SOFTMAX_ACCUMULATOR_VAL_ZERO : exp_sums[exp_sums_offset + idx * SUBGROUP_SIZE + sglid]; + } + + SOFTMAX_ACCUMULATOR_TYPE global_max_logit = max_logit[0]; + for (uint i = 1; i < CEIL_DIV(partitions_num, SUBGROUP_SIZE); i++) { + global_max_logit = SOFTMAX_ACCUMULATOR_MAX_FUNC(global_max_logit, max_logit[i]); + } + + global_max_logit = sub_group_reduce_max(global_max_logit); + + SOFTMAX_ACCUMULATOR_TYPE global_exp_sum = SOFTMAX_ACCUMULATOR_VAL_ZERO; + for (uint i = 0; i < CEIL_DIV(partitions_num, SUBGROUP_SIZE); i++) { + SOFTMAX_ACCUMULATOR_TYPE adjusted_exp_sum = exp_sum[i] * native_exp(max_logit[i] - global_max_logit); + // slm_exp_sums[head_idx][i * SUBGROUP_SIZE + sglid] = adjusted_exp_sum; + if (i * SUBGROUP_SIZE + sglid == partition_idx) + slm_exp_sums[head_idx] = adjusted_exp_sum; + global_exp_sum += adjusted_exp_sum; + } + + global_exp_sum = sub_group_reduce_add(global_exp_sum); + + slm_global_exp_sum[head_idx] = global_exp_sum; + } + + barrier(CLK_LOCAL_MEM_FENCE); + + for (uint head_idx = 0; head_idx < HEADS_NUM; head_idx++) { + SOFTMAX_ACCUMULATOR_TYPE adjusted_exp_sum = slm_exp_sums[head_idx]; + SOFTMAX_ACCUMULATOR_TYPE global_exp_sum = slm_global_exp_sum[head_idx]; + + const uint input_offset = subsequence_idx * HEADS_NUM * max_seq_len + head_idx * max_seq_len + partition_global_idx; + SOFTMAX_ACCUMULATOR_TYPE softmax_value = softmax_output[input_offset]; + + softmax_value = softmax_value * adjusted_exp_sum / global_exp_sum; + total_score += softmax_value; + } + } else { + // Non optimized fallback version + const uint subsequence_pos = is_mixed_mode ? subsequence_end - 1 : subsequence_idx; + for (uint head_idx = 0; head_idx < HEADS_NUM; head_idx++) { + SOFTMAX_ACCUMULATOR_TYPE global_max_logit = SOFTMAX_ACCUMULATOR_VAL_MIN; + const uint max_logits_base_offset = subsequence_pos * HEADS_NUM * partitions_num + head_idx * partitions_num; + for (uint i = 0; i < CEIL_DIV(partitions_num, SUBGROUP_SIZE); i++) { + const uint partition_offset = i * SUBGROUP_SIZE + sglid; + SOFTMAX_ACCUMULATOR_TYPE max_logit = partition_offset >= partitions_num ? SOFTMAX_ACCUMULATOR_VAL_MIN : max_logits[max_logits_base_offset + partition_offset]; + global_max_logit = SOFTMAX_ACCUMULATOR_MAX_FUNC(global_max_logit, max_logit); + } + + global_max_logit = sub_group_reduce_max(global_max_logit); + + SOFTMAX_ACCUMULATOR_TYPE global_exp_sum = SOFTMAX_ACCUMULATOR_VAL_ZERO; + SOFTMAX_ACCUMULATOR_TYPE partition_adjusted_exp_sum = SOFTMAX_ACCUMULATOR_VAL_ZERO; + const uint exp_sums_base_offset = subsequence_pos * HEADS_NUM * partitions_num + head_idx * partitions_num; + for (uint i = 0; i < CEIL_DIV(partitions_num, SUBGROUP_SIZE); i++) { + const uint partition_offset = i * SUBGROUP_SIZE + sglid; + SOFTMAX_ACCUMULATOR_TYPE exp_sum = partition_offset >= partitions_num ? SOFTMAX_ACCUMULATOR_VAL_ZERO : exp_sums[exp_sums_base_offset + partition_offset]; + SOFTMAX_ACCUMULATOR_TYPE max_logit = partition_offset >= partitions_num ? SOFTMAX_ACCUMULATOR_VAL_MIN : max_logits[max_logits_base_offset + partition_offset]; + SOFTMAX_ACCUMULATOR_TYPE adjusted_exp_sum = exp_sum * native_exp(max_logit - global_max_logit); + global_exp_sum += adjusted_exp_sum; + + // Save and broadcast the adjusted exp_sum for the currently being processed partition + if (i == partition_idx / SUBGROUP_SIZE) + partition_adjusted_exp_sum = sub_group_broadcast(adjusted_exp_sum, partition_idx % SUBGROUP_SIZE); + } + + global_exp_sum = sub_group_reduce_add(global_exp_sum); + + const uint input_offset = subsequence_idx * HEADS_NUM * max_seq_len + head_idx * max_seq_len + partition_global_idx; + SOFTMAX_ACCUMULATOR_TYPE softmax_value = softmax_output[input_offset]; + + softmax_value = softmax_value * partition_adjusted_exp_sum / global_exp_sum; + total_score += softmax_value; + } + } + + const uint output_offset = subsequence_offsets[subsequence_idx]; + if (partition_global_idx < seq_len) { + scores_output[output_offset + partition_global_idx] = total_score; + } +} + +#undef MAX_PARTITIONS_NUM +#endif diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data_b_fs_yx_fsv16_fsv32_to_bfyx.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data_b_fs_yx_fsv16_fsv32_to_bfyx.cl index 95f0d0ff399a3b..ee27d220e30ce9 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data_b_fs_yx_fsv16_fsv32_to_bfyx.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data_b_fs_yx_fsv16_fsv32_to_bfyx.cl @@ -66,10 +66,7 @@ KERNEL (reorder_data_b_fs_yx_fsv16_fsv32_to_bfyx)( #if (TILE_SIZE == DEFAULT_TILE_SIZE) - // read - INPUTVTYPE read_data = AS_INPUTVTYPE(_sub_group_block_read8((const __global uint*)(input) + input_idx_tile)); - - // write + // write index const uint output_idx = OUTPUT_GET_TILED_INDEX(OUTPUT_TILED_ORDER); if (F_NO_REMAINDER_CONDITION @@ -79,13 +76,25 @@ KERNEL (reorder_data_b_fs_yx_fsv16_fsv32_to_bfyx)( ) { #ifdef X_REMAINDER_SIZE if (X_REMAINDER_CONDITION) { + // read + INPUTVTYPE read_data; + for (int j = 0; j < X_REMAINDER_SIZE; ++j) { + read_data[j] = AS_INPUT0_TYPE(_sub_group_block_read((const __global uint*)(input) + input_idx_tile + j * DEFAULT_STRIDE)); + } + // write for (int i = 0 ; i < X_REMAINDER_SIZE; i++) { output[output_idx + i] = TO_OUTPUT_TYPE(read_data[i]); } } else { + // read + INPUTVTYPE read_data = AS_INPUTVTYPE(_sub_group_block_read8((const __global uint*)(input) + input_idx_tile)); + // write VSTORE(TO_OUTPUTVTYPE(read_data), 0, output + output_idx); } #else + // read + INPUTVTYPE read_data = AS_INPUTVTYPE(_sub_group_block_read8((const __global uint*)(input) + input_idx_tile)); + // write VSTORE(TO_OUTPUTVTYPE(read_data), 0, output + output_idx); #endif } diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data_bfyx_to_blocked_format.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data_bfyx_to_blocked_format.cl index 45d0ccc5c0933e..2f403b798dea39 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data_bfyx_to_blocked_format.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data_bfyx_to_blocked_format.cl @@ -26,6 +26,18 @@ } \ } +#define FUNC_LOAD_LEFTOVERS(inner, outer) unroll_for (uint lh = 0; lh < outer; ++lh) { \ + const uint input_idx = INPUT0_GET_TILED_INDEX(INPUT0_TILED_ORDER); \ + INPUTVTYPE read_data; \ + unroll_for (uint lw = 0; lw < inner; ++lw) { \ + read_data[lw] = input[input_idx + lw]; \ + } \ + unroll_for (uint lw = 0; lw < inner; ++lw) { \ + const uint dst = local_buf_offset + lw; \ + transpose_buf[dst][lh] = read_data[lw]; \ + } \ + } + #define FUNC_VSTORE(loop) unroll_for (uint lw = 0; lw < loop; ++lw) { \ const uint output_idx = output_idx_tile + (lw * x_pitch); \ VSTORE(TO_OUTPUTVTYPE(transpose_buf[local_buf_offset + lw]), 0, output + output_idx); \ @@ -109,7 +121,15 @@ KERNEL (reorder_data_bfyx_to_blocked_format)( if (F_NO_REMAINDER_CONDITION) { // read and transpose +#ifdef X_REMAINDER_CONDITION + if (X_NO_REMAINDER_CONDITION) { + FUNC_VLOAD(TILE_SIZE, TILE_SIZE) + } else { + FUNC_LOAD_LEFTOVERS(X_REMAINDER_SIZE, TILE_SIZE) + } +#else FUNC_VLOAD(TILE_SIZE, TILE_SIZE) +#endif // write to ddr #ifdef X_REMAINDER_CONDITION @@ -125,7 +145,15 @@ KERNEL (reorder_data_bfyx_to_blocked_format)( #ifdef F_REMAINDER_CONDITION else if (F_REMAINDER_CONDITION) { // read and transpose + #ifdef X_REMAINDER_CONDITION + if (X_NO_REMAINDER_CONDITION) { + FUNC_VLOAD(TILE_SIZE, F_REMAINDER_SIZE) + } else { + FUNC_LOAD_LEFTOVERS(X_REMAINDER_SIZE, F_REMAINDER_SIZE) + } + #else FUNC_VLOAD(TILE_SIZE, F_REMAINDER_SIZE) + #endif // write to ddr #ifdef X_REMAINDER_CONDITION diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/sdpa_opt.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/sdpa_opt.cl index 55f87e4189d9fe..cddafe62623d9e 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/sdpa_opt.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/sdpa_opt.cl @@ -842,6 +842,14 @@ KERNEL(sdpa_opt)( const __global int* blocked_indexes_start, const __global int* blocked_indexes_end, const __global int* gws_seq_indexes_correspondence +#if PAGED_ATTENTION_SCORES_OUTPUT + , __global SOFTMAX_ACCUMULATOR_TYPE* softmax_results + , const __global int* subsequence_offsets + , __global SOFTMAX_ACCUMULATOR_TYPE* exp_sums + , __global SOFTMAX_ACCUMULATOR_TYPE* max_logits + , __global OUTPUT_TYPE* tmp_out + , const uint aligned_max_context_len +#endif #else __global SOFTMAX_ACCUMULATOR_TYPE* exp_sums, __global SOFTMAX_ACCUMULATOR_TYPE* max_logits, @@ -1222,6 +1230,39 @@ KERNEL(sdpa_opt)( slm_qk_vals[sglid * SEQ_LEN_PARTITION_SIZE + sgid * TARGET_SEQ_LEN_BLOCK_SIZE + i] = qk_acc[i]; } +#if PAGED_ATTENTION_SCORES_OUTPUT + const uint subsequence_idx = gws_seq_indexes_correspondence[target_seq_dim]; + const uint subsequence_end_pos = subsequence_begins[subsequence_idx + 1]; + const uint block_start_pos = blocked_indexes_start[target_seq_dim]; + const uint block_end_pos = blocked_indexes_end[target_seq_dim]; + + // PagedAttention is supposed to save only last "row" of the QK matrix multiplication, + // so save SEQ_LEN_PARTITION_SIZE elements for each partition + if (subsequence_end_pos == block_end_pos) { + const uint last_row_idx = block_end_pos - block_start_pos - 1; + if (sglid == last_row_idx) { + const uint partition_idx = start_partition_idx / SEQ_LEN_PARTITION_SIZE; + + if (sgid == 0) { + const uint max_partitions_num = aligned_max_context_len / SEQ_LEN_PARTITION_SIZE; + const uint exp_sums_output_offset = subsequence_idx * NUM_HEADS * max_partitions_num + + num_heads_dim * max_partitions_num + + partition_idx; + exp_sums[exp_sums_output_offset] = exp_sum_new; + max_logits[exp_sums_output_offset] = qk_max_new; + } + + const uint output_offset = subsequence_idx * NUM_HEADS * aligned_max_context_len + + num_heads_dim * aligned_max_context_len + + partition_idx * SEQ_LEN_PARTITION_SIZE + sgid * TARGET_SEQ_LEN_BLOCK_SIZE; + for (uint i = 0; i < TARGET_SEQ_LEN_BLOCK_SIZE; i++) { + softmax_results[output_offset + i] = qk_acc[i]; + } + + } + } +#endif + barrier(CLK_LOCAL_MEM_FENCE); } diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/pa_kv_cache_update_kernel_ref.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/pa_kv_cache_update_kernel_ref.cpp index ddfb491f50278a..ce20f49de597ff 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/pa_kv_cache_update_kernel_ref.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/pa_kv_cache_update_kernel_ref.cpp @@ -167,7 +167,7 @@ void KVCacheUpdateKernelRef::GetUpdateDispatchDataFunc(KernelData& kd) const { const auto indexes_dt = Datatype::INT32; const auto target_seq_len_block_size = 16; - const auto target_seq_len = prim_params.conf.paged_attention_aligned_seq_len; + const auto target_seq_len = std::max(prim_params.conf.paged_attention_aligned_seq_len, static_cast(1)); const auto indexes_buf_size = CeilDiv(target_seq_len, target_seq_len_block_size) * BytesPerElement(indexes_dt); kd.internalBufferSizes.clear(); diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/pa_sdpa_kernel_opt.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/pa_sdpa_kernel_opt.cpp index 63c5e74160f652..909a40d677f535 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/pa_sdpa_kernel_opt.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/pa_sdpa_kernel_opt.cpp @@ -2,6 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 // +#include "sdpa_kernel_opt.h" #include "pa_sdpa_kernel_opt.h" #include "kernel_selector_params.h" @@ -15,6 +16,7 @@ enum KernelsTypes { MULTI_TOKENS, FINALIZATION, FINALIZATION_MULTI_TOKENS, + SCORES_CALCULATION, TOTAL_KERNELS_NUM }; @@ -35,6 +37,8 @@ static std::string GetKernelName(std::string base_name, KernelsTypes type) { kernel_name += "_finalization"; } else if (type == KernelsTypes::FINALIZATION_MULTI_TOKENS) { kernel_name += "_finalization_multi_tokens_seq"; + } else if (type == KernelsTypes::SCORES_CALCULATION) { + kernel_name += "_scores_calculation"; } return kernel_name; @@ -46,10 +50,15 @@ KernelsData PagedAttentionSDPAKernelOpt::GetKernelsData(const Params& p) const { } const auto& params = static_cast(p); - const std::vector kernels_type = { KernelsTypes::SINGLE_TOKEN, - KernelsTypes::MULTI_TOKENS, - KernelsTypes::FINALIZATION, - KernelsTypes::FINALIZATION_MULTI_TOKENS }; + std::vector kernels_type = { KernelsTypes::SINGLE_TOKEN, + KernelsTypes::MULTI_TOKENS, + KernelsTypes::FINALIZATION, + KernelsTypes::FINALIZATION_MULTI_TOKENS }; + + const auto has_scores_output = params.outputs.size() > 1; + if (has_scores_output) { + kernels_type.push_back(KernelsTypes::SCORES_CALCULATION); + } KernelData kd = KernelData::Default(params, kernels_type.size()); kd.needs_sub_kernels_sync = true; @@ -65,7 +74,8 @@ KernelsData PagedAttentionSDPAKernelOpt::GetKernelsData(const Params& p) const { const auto jit = CreateJit(kernel_name, jit_constants, entry_point); - size_t inputs_num = static_cast(params.inputs.size()); + int inputs_num = static_cast(params.inputs.size()); + int outputs_num = 1; if (kernel_type == KernelsTypes::SINGLE_TOKEN) { // SINGLE_TOKEN kernel doesn't use the subsequence_begins input inputs_num -= 1; @@ -75,6 +85,11 @@ KernelsData PagedAttentionSDPAKernelOpt::GetKernelsData(const Params& p) const { } else if (kernel_type == KernelsTypes::FINALIZATION_MULTI_TOKENS) { // FINALIZATION_MULTI_TOKENS kernel uses past_lens data input and subsequence_begins inputs_num = 2; + } else if (kernel_type == KernelsTypes::SCORES_CALCULATION) { + // SCORES_CALCULATION kernel uses past_lens data input and subsequence_begins + inputs_num = 2; + // Output is configured manually to use the second output memory buffer + outputs_num = 0; } auto& kernel = kd.kernels[kd_kernels_idx++]; @@ -87,19 +102,33 @@ KernelsData PagedAttentionSDPAKernelOpt::GetKernelsData(const Params& p) const { {}, false, false, - static_cast(inputs_num), + inputs_num, GetFusedPrimitiveInputsCount(params), - static_cast(params.outputs.size()), + outputs_num, params.is_shape_agnostic); - 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}); + if (kernel_type == KernelsTypes::SCORES_CALCULATION) { + kernel.params.arguments.push_back({ArgumentDescriptor::Types::OUTPUT, 1}); + } + + uint32_t internal_buffers_num = 0; + if (has_scores_output) { + // Intermediate softmax results for scores output calculation and precalculated accumulated + // sequence length offsets for each subsequence + internal_buffers_num += 2; + } + + // Softmax's exp_sums, max_logits and intermediate output + internal_buffers_num += 3; if (kernel_type == KernelsTypes::MULTI_TOKENS || kernel_type == KernelsTypes::FINALIZATION_MULTI_TOKENS) { // MULTIPLE_TOKENS kernels needs additional information related to mapping // launched kernel instances to subsequence indexes - kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 3}); + internal_buffers_num++; + } + + for (uint32_t i = 0; i < internal_buffers_num; i++) { + kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, i}); } if (kernel_type == KernelsTypes::FINALIZATION || kernel_type == KernelsTypes::FINALIZATION_MULTI_TOKENS) { @@ -108,6 +137,15 @@ KernelsData PagedAttentionSDPAKernelOpt::GetKernelsData(const Params& p) const { // Remove unused shape_info argument at finalization stage kernel.params.arguments.erase(kernel.params.arguments.begin()); } + + if (kernel_type == KernelsTypes::SCORES_CALCULATION) { + // The scores kernel needs to know if the current execution mode is mixed or ordinary + // to configure proper memory access + kernel.params.arguments.push_back({ArgumentDescriptor::Types::SCALAR, 0}); + + // Remove unused shape_info argument for scores kernel + kernel.params.arguments.erase(kernel.params.arguments.begin()); + } } return {kd}; @@ -173,7 +211,12 @@ JitConstants PagedAttentionSDPAKernelOpt::GetJitConstants(const pa_sdpa_params& jit.AddConstant(MakeJitConstant("BROADCAST_GROUP_SIZE", config.group_size)); } - auto sdpa_stage = kernel_idx == KernelsTypes::FINALIZATION || kernel_idx == KernelsTypes::FINALIZATION_MULTI_TOKENS ? 1 : 0; + auto sdpa_stage = 0; + if (kernel_idx == KernelsTypes::FINALIZATION || kernel_idx == KernelsTypes::FINALIZATION_MULTI_TOKENS) { + sdpa_stage = 1; + } else if (kernel_idx == KernelsTypes::SCORES_CALCULATION) { + sdpa_stage = 2; + } jit.AddConstant(MakeJitConstant("SDPA_STAGE_" + std::to_string(sdpa_stage), 1)); if (config.has_const_scale_val) { @@ -190,6 +233,10 @@ JitConstants PagedAttentionSDPAKernelOpt::GetJitConstants(const pa_sdpa_params& jit.Merge(MakeTypeJitConstants(params.inputs[alibi_input_idx].GetDType(), "ALIBI_INPUT")); } + if (params.outputs.size() > 1) { + jit.AddConstant(MakeJitConstant("PAGED_ATTENTION_SCORES_OUTPUT", 1)); + } + if (kernel_idx == KernelsTypes::MULTI_TOKENS || kernel_idx == KernelsTypes::FINALIZATION_MULTI_TOKENS) jit.AddConstant(MakeJitConstant("MULTI_TOKENS_PROCESSING", 1)); @@ -203,18 +250,36 @@ CommonDispatchData PagedAttentionSDPAKernelOpt::SetDefault(const pa_sdpa_params& const auto& input = params.inputs[0]; if (!input.is_dynamic()) { - const size_t sequences_number = input.Batch().v; - const size_t num_of_partitions = CeilDiv(params.max_context_len, seq_len_partition_size); + const size_t total_tokens = input.Batch().v; + const size_t num_of_partitions = CeilDiv(params.conf.paged_attention_max_len, seq_len_partition_size); const size_t heads_num = static_cast(params.conf.heads_num); const size_t head_size = static_cast(params.conf.head_size); - if (kernel_idx == 0) { - dispatch_data.gws = { sequences_number, + if (kernel_idx == KernelsTypes::SINGLE_TOKEN || kernel_idx == KernelsTypes::MULTI_TOKENS) { + dispatch_data.gws = { total_tokens, heads_num, head_size * num_of_partitions }; dispatch_data.lws = { 1, 1, head_size }; + } else if (kernel_idx == KernelsTypes::SCORES_CALCULATION) { + const auto& past_lens = params.inputs[3]; + const auto subsequences_number = past_lens.Batch().v; + + size_t partition_size = 0; + size_t num_of_partitions = 0; + if (params.stage == PagedAttentionStage::PREFILL) { + partition_size = SDPAKernelOpt::get_seq_len_partition_size(params, params.conf.head_size, 1); + } else { + partition_size = seq_len_partition_size; + } + + num_of_partitions = CeilDiv(params.conf.paged_attention_max_len, partition_size); + + dispatch_data.gws = { partition_size * num_of_partitions, + 1, + subsequences_number }; + dispatch_data.lws = { partition_size, 1, 1 }; } else { - dispatch_data.gws = { sequences_number, + dispatch_data.gws = { total_tokens, heads_num, head_size }; dispatch_data.lws = { 1, 1, subgroup_size }; @@ -228,30 +293,39 @@ void PagedAttentionSDPAKernelOpt::GetUpdateDispatchDataFunc(KernelData& kd) cons kd.update_dispatch_data_func = [](const Params& params, KernelData& kd) { const auto& prim_params = static_cast(params); - const size_t expected_kernels_num = 4; - OPENVINO_ASSERT(kd.kernels.size() == expected_kernels_num, "[GPU] Invalid kernels size for update dispatch data func of SDPA kernel"); + const auto has_scores_output = prim_params.outputs.size() > 1; + const auto expected_kernels_num = has_scores_output ? KernelsTypes::TOTAL_KERNELS_NUM : KernelsTypes::TOTAL_KERNELS_NUM - 1; + OPENVINO_ASSERT(kd.kernels.size() == static_cast(expected_kernels_num), + "[GPU] Invalid kernels size for update dispatch data func of SDPA kernel"); + + const auto scores_calc_only = prim_params.stage == PagedAttentionStage::PREFILL && has_scores_output; + const auto multi_tokens_mode = prim_params.stage == PagedAttentionStage::MIXED; auto dispatch_data1 = SetDefault(prim_params, KernelsTypes::SINGLE_TOKEN); kd.kernels[KernelsTypes::SINGLE_TOKEN].params.workGroups.global = dispatch_data1.gws; kd.kernels[KernelsTypes::SINGLE_TOKEN].params.workGroups.local = dispatch_data1.lws; - kd.kernels[KernelsTypes::SINGLE_TOKEN].skip_execution = prim_params.multi_tokens_mode; + kd.kernels[KernelsTypes::SINGLE_TOKEN].skip_execution = multi_tokens_mode || scores_calc_only; kd.kernels[KernelsTypes::MULTI_TOKENS].params.workGroups.global = dispatch_data1.gws; kd.kernels[KernelsTypes::MULTI_TOKENS].params.workGroups.local = dispatch_data1.lws; - kd.kernels[KernelsTypes::MULTI_TOKENS].skip_execution = !prim_params.multi_tokens_mode; + kd.kernels[KernelsTypes::MULTI_TOKENS].skip_execution = !multi_tokens_mode || scores_calc_only; - const auto& input = prim_params.inputs[0]; - const size_t sequences_number = input.Batch().v; - const size_t num_of_partitions = CeilDiv(prim_params.max_context_len, seq_len_partition_size); + size_t partition_size = 0; + if (prim_params.stage == PagedAttentionStage::PREFILL) { + partition_size = SDPAKernelOpt::get_seq_len_partition_size(params, prim_params.conf.head_size, 1); + } else { + partition_size = seq_len_partition_size; + } + const size_t num_of_partitions = CeilDiv(prim_params.conf.paged_attention_max_len, partition_size); auto dispatch_data2 = SetDefault(prim_params, KernelsTypes::FINALIZATION); kd.kernels[KernelsTypes::FINALIZATION].params.workGroups.global = dispatch_data2.gws; kd.kernels[KernelsTypes::FINALIZATION].params.workGroups.local = dispatch_data2.lws; - kd.kernels[KernelsTypes::FINALIZATION].skip_execution = num_of_partitions == 1 || prim_params.multi_tokens_mode; + kd.kernels[KernelsTypes::FINALIZATION].skip_execution = num_of_partitions == 1 || multi_tokens_mode || scores_calc_only; kd.kernels[KernelsTypes::FINALIZATION_MULTI_TOKENS].params.workGroups.global = dispatch_data2.gws; kd.kernels[KernelsTypes::FINALIZATION_MULTI_TOKENS].params.workGroups.local = dispatch_data2.lws; - kd.kernels[KernelsTypes::FINALIZATION_MULTI_TOKENS].skip_execution = num_of_partitions == 1 || !prim_params.multi_tokens_mode; + kd.kernels[KernelsTypes::FINALIZATION_MULTI_TOKENS].skip_execution = num_of_partitions == 1 || !multi_tokens_mode || scores_calc_only; ScalarDescriptor num_of_partitions_scalar; num_of_partitions_scalar.t = ScalarDescriptor::Types::UINT32; @@ -261,23 +335,63 @@ void PagedAttentionSDPAKernelOpt::GetUpdateDispatchDataFunc(KernelData& kd) cons kd.kernels[KernelsTypes::FINALIZATION_MULTI_TOKENS].params.scalars.resize(1); kd.kernels[KernelsTypes::FINALIZATION_MULTI_TOKENS].params.scalars[0] = num_of_partitions_scalar; + if (has_scores_output) { + auto dispatch_data = SetDefault(prim_params, KernelsTypes::SCORES_CALCULATION); + kd.kernels[KernelsTypes::SCORES_CALCULATION].params.workGroups.global = dispatch_data.gws; + kd.kernels[KernelsTypes::SCORES_CALCULATION].params.workGroups.local = dispatch_data.lws; + kd.kernels[KernelsTypes::SCORES_CALCULATION].skip_execution = false; + + ScalarDescriptor is_mixed_mode; + is_mixed_mode.t = ScalarDescriptor::Types::UINT32; + is_mixed_mode.v.u32 = static_cast(multi_tokens_mode); + kd.kernels[KernelsTypes::SCORES_CALCULATION].params.scalars.resize(1); + kd.kernels[KernelsTypes::SCORES_CALCULATION].params.scalars[0] = is_mixed_mode; + } + + const auto& input = prim_params.inputs[0]; + const size_t total_tokens = input.Batch().v; + auto buf_dt_size = BytesPerElement(softmax_acc_dt); - auto buf_elements_count = sequences_number * prim_params.conf.heads_num * num_of_partitions; + auto buf_elements_count = total_tokens * prim_params.conf.heads_num * num_of_partitions; auto buf_size = buf_elements_count * buf_dt_size; auto tmp_out_dt_size = BytesPerElement(softmax_acc_dt); - auto tmp_out_elements_count = sequences_number * prim_params.conf.heads_num * prim_params.conf.head_size * num_of_partitions; + auto tmp_out_elements_count = total_tokens * prim_params.conf.heads_num * prim_params.conf.head_size * num_of_partitions; auto tmp_out_size = tmp_out_elements_count * tmp_out_dt_size; kd.internalBufferSizes.clear(); - kd.internalBufferSizes.push_back(buf_size); - kd.internalBufferSizes.push_back(buf_size); - kd.internalBufferSizes.push_back(tmp_out_size); + + if (has_scores_output) { + const auto& past_lens = prim_params.inputs[3]; + auto subsequences_number = past_lens.Batch().v; + auto softmax_buf_dt_size = BytesPerElement(softmax_acc_dt); + + auto softmax_buf_elements_count = subsequences_number * prim_params.conf.heads_num * num_of_partitions * partition_size; + auto softmax_buf_size = softmax_buf_elements_count * softmax_buf_dt_size; + + // Softmax intermediate output + kd.internalBufferSizes.push_back(softmax_buf_size); + // Precalculated accumulated sequence length offsets for each subsequence + kd.internalBufferSizes.push_back(subsequences_number * BytesPerElement(Datatype::INT32)); + + if (prim_params.stage == PagedAttentionStage::PREFILL) { + // Recalculate buf_size as in case of PREFILL stage it's not needed to allocate buffer per each input token + buf_elements_count = subsequences_number * prim_params.conf.heads_num * num_of_partitions; + buf_size = buf_elements_count * buf_dt_size; + + // Intermediate tmp output buffer is not used for PREFILL stage + tmp_out_size = tmp_out_dt_size; + } + } + + kd.internalBufferSizes.push_back(buf_size); // softmax exp_sums + kd.internalBufferSizes.push_back(buf_size); // softmax max_logits + kd.internalBufferSizes.push_back(tmp_out_size); // intermediate output kd.internalBufferDataType = softmax_acc_dt; - if (prim_params.multi_tokens_mode) { + if (multi_tokens_mode) { auto buf_dt_size = BytesPerElement(Datatype::INT32); - auto buf_elements_count = sequences_number; + auto buf_elements_count = total_tokens; auto buf_size = Align(buf_elements_count * buf_dt_size, BytesPerElement(softmax_acc_dt)); kd.internalBufferSizes.push_back(buf_size); } diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/pa_sdpa_kernel_opt.h b/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/pa_sdpa_kernel_opt.h index a2456ccd9e2af5..a52571b03691df 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/pa_sdpa_kernel_opt.h +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/pa_sdpa_kernel_opt.h @@ -9,11 +9,17 @@ namespace kernel_selector { +enum PagedAttentionStage { + GENERATE = 0, + PREFILL = 1, + MIXED = 2, + UNKNOWN = 3 +}; + struct pa_sdpa_params : base_params { pa_sdpa_params() : base_params(KernelType::PA_SDPA) {} - bool multi_tokens_mode = false; - size_t max_context_len = 0; + PagedAttentionStage stage = PagedAttentionStage::UNKNOWN; sdpa_configuration conf; }; diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_base.h b/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_base.h index 5cd9c384ff2709..8fcc4a16692d6c 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_base.h +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_base.h @@ -97,6 +97,7 @@ struct sdpa_configuration { bool is_paged_attention = false; int64_t paged_attention_aligned_seq_len = -1; int64_t paged_attention_block_size = 0; + int64_t paged_attention_max_len = 0; bool has_const_scale_val = false; float scale_val = 0.f; }; diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_opt.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_opt.cpp index 4e71064efbc895..4c23d4de4fd68d 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_opt.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_opt.cpp @@ -21,38 +21,11 @@ enum KernelsTypes { constexpr size_t subgroup_size = 16; } // namespace -static size_t get_sg_number_scale_factor(const sdpa_params& sdpa_params, size_t kernel_type) { - const size_t optimal_scale_factor = 2; - if (kernel_type == KernelsTypes::MULTI_TOKENS) { - if (sdpa_params.conf.head_size * optimal_scale_factor <= sdpa_params.engineInfo.maxWorkGroupSize) { - return optimal_scale_factor; - } - } else if (kernel_type == KernelsTypes::SINGLE_TOKEN) { - if (sdpa_params.conf.head_size * optimal_scale_factor <= sdpa_params.engineInfo.maxWorkGroupSize && - sdpa_params.conf.head_size * optimal_scale_factor / subgroup_size <= subgroup_size) { - return optimal_scale_factor; - } - } - - return 1; -} - static size_t get_target_seq_len_block_size() { const size_t block_size = 16; return block_size; } -static size_t get_seq_len_partition_size(const sdpa_params& sdpa_params, size_t kernel_type) { - size_t seq_len = 0; - if (kernel_type == KernelsTypes::MULTI_TOKENS) { - seq_len = sdpa_params.conf.head_size * get_sg_number_scale_factor(sdpa_params, kernel_type); - } else { - seq_len = 256; - } - - return seq_len; -} - static Datatype get_softmax_acc_type() { return Datatype::F32; } @@ -71,7 +44,7 @@ static size_t get_partitions_num(const sdpa_params& sdpa_params, size_t kernel_t TransposedDimensionAccessHelperBase dims_k(sdpa_params.inputs[1], sdpa_params.input1_order); auto source_seq_len = dims_k.y_dim().v; - return CeilDiv(source_seq_len, get_seq_len_partition_size(sdpa_params, kernel_type)); + return CeilDiv(source_seq_len, SDPAKernelOpt::get_seq_len_partition_size(sdpa_params, sdpa_params.conf.head_size, kernel_type)); } static std::vector get_internal_buffer_sizes(const sdpa_params& sdpa_params, size_t kernel_type) { @@ -130,6 +103,33 @@ static std::string GetKernelName(std::string base_name, KernelsTypes type, const return kernel_name; } +size_t SDPAKernelOpt::get_sg_number_scale_factor(const Params& params, size_t head_size, size_t kernel_type) { + const size_t optimal_scale_factor = 2; + if (kernel_type == KernelsTypes::MULTI_TOKENS) { + if (head_size * optimal_scale_factor <= params.engineInfo.maxWorkGroupSize) { + return optimal_scale_factor; + } + } else if (kernel_type == KernelsTypes::SINGLE_TOKEN) { + if (head_size * optimal_scale_factor <= params.engineInfo.maxWorkGroupSize && + head_size * optimal_scale_factor / subgroup_size <= subgroup_size) { + return optimal_scale_factor; + } + } + + return 1; +} + +size_t SDPAKernelOpt::get_seq_len_partition_size(const Params& params, size_t head_size, size_t kernel_type) { + size_t seq_len = 0; + if (kernel_type == KernelsTypes::MULTI_TOKENS) { + seq_len = head_size * get_sg_number_scale_factor(params, head_size, kernel_type); + } else { + seq_len = 256; + } + + return seq_len; +} + ParamsKey SDPAKernelOpt::GetSupportedKey() const { ParamsKey k; k.EnableInputDataType(Datatype::INT8); @@ -176,14 +176,14 @@ JitConstants SDPAKernelOpt::GetJitConstants(const sdpa_params& params, size_t ke const auto& config = params.conf; jit.AddConstant(MakeJitConstant("SUBGROUP_SIZE", subgroup_size)); jit.AddConstant(MakeJitConstant("HEAD_SIZE", config.head_size)); - jit.AddConstant(MakeJitConstant("SEQ_LEN_PARTITION_SIZE", get_seq_len_partition_size(params, kernel_idx))); + jit.AddConstant(MakeJitConstant("SEQ_LEN_PARTITION_SIZE", get_seq_len_partition_size(params, config.head_size, kernel_idx))); auto target_seq_len_block_size = kernel_idx == KernelsTypes::SINGLE_TOKEN ? 1 : get_target_seq_len_block_size(); jit.AddConstant(MakeJitConstant("TARGET_SEQ_LEN_BLOCK_SIZE", target_seq_len_block_size)); auto sdpa_stage = kernel_idx == KernelsTypes::FINALIZATION ? 1 : 0; jit.AddConstant(MakeJitConstant("SDPA_STAGE_" + std::to_string(sdpa_stage), 1)); - jit.AddConstant(MakeJitConstant("SG_SCALE_FACTOR", get_sg_number_scale_factor(params, kernel_idx))); + jit.AddConstant(MakeJitConstant("SG_SCALE_FACTOR", get_sg_number_scale_factor(params, config.head_size, kernel_idx))); if (params.conf.is_paged_attention) { if (params.conf.has_alibi_input) { @@ -196,6 +196,10 @@ JitConstants SDPAKernelOpt::GetJitConstants(const sdpa_params& params, size_t ke } else { jit.AddConstant(MakeJitConstant("HAS_SCALE_INPUT", 1)); } + + if (params.outputs.size() > 1) { + jit.AddConstant(MakeJitConstant("PAGED_ATTENTION_SCORES_OUTPUT", 1)); + } } else if (params.inputs.size() <= 4) { jit.AddConstant(MakeJitConstant("STATIC_SCALE_VALUE_INV", std::sqrt(static_cast(params.conf.head_size)))); jit.AddConstant(MakeJitConstant("STATIC_SCALE_VALUE", 1.0f / std::sqrt(static_cast(params.conf.head_size)))); @@ -218,11 +222,11 @@ CommonDispatchData SDPAKernelOpt::SetDefault(const sdpa_params& params, size_t k if (params.conf.is_paged_attention) { OPENVINO_ASSERT(kernel_idx == KernelsTypes::MULTI_TOKENS); - const size_t sg_num_scale = get_sg_number_scale_factor(params, kernel_idx); const size_t heads_num = static_cast(params.conf.heads_num); + const size_t head_size = static_cast(params.conf.head_size); + const size_t sg_num_scale = get_sg_number_scale_factor(params, head_size, kernel_idx); const size_t target_seq_len_block_size = get_target_seq_len_block_size(); const size_t target_seq_len = static_cast(params.conf.paged_attention_aligned_seq_len); - const size_t head_size = static_cast(params.conf.head_size); dispatch_data.gws = { heads_num, CeilDiv(target_seq_len, target_seq_len_block_size), @@ -243,13 +247,13 @@ CommonDispatchData SDPAKernelOpt::SetDefault(const sdpa_params& params, size_t k const size_t target_seq_len_block_size = kernel_idx == 1 ? get_target_seq_len_block_size() : 1; if (kernel_idx == KernelsTypes::SINGLE_TOKEN) { - const size_t sg_num_scale = get_sg_number_scale_factor(params, kernel_idx); + const size_t sg_num_scale = get_sg_number_scale_factor(params, head_size, kernel_idx); dispatch_data.gws = { batch_size * heads_num, CeilDiv(target_seq_len, target_seq_len_block_size), head_size * num_of_partitions * sg_num_scale }; dispatch_data.lws = { 1, 1, head_size * sg_num_scale }; } else if (kernel_idx == KernelsTypes::MULTI_TOKENS) { - const size_t sg_num_scale = get_sg_number_scale_factor(params, kernel_idx); + const size_t sg_num_scale = get_sg_number_scale_factor(params, head_size, kernel_idx); dispatch_data.gws = { batch_size * heads_num, CeilDiv(target_seq_len, target_seq_len_block_size), head_size * sg_num_scale }; @@ -317,7 +321,7 @@ KernelsData SDPAKernelOpt::GetKernelsData(const Params& params) const { false, inputs_num, GetFusedPrimitiveInputsCount(params), - static_cast(prim_params.outputs.size()), + 1 /* number_of_outputs */, prim_params.is_shape_agnostic); auto beam_table_idx = prim_params.inputs.size(); @@ -339,6 +343,19 @@ KernelsData SDPAKernelOpt::GetKernelsData(const Params& params) const { kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 1}); kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 2}); + if (prim_params.conf.is_paged_attention && prim_params.outputs.size() > 1) { + // Intermediate buffers for PagedAttention scores calculation: + // softmax_results, subsequence_offsets, exp_sums, max_logits, tmp_out + kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 3}); + kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 4}); + kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 5}); + kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 6}); + kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 7}); + + // Scalar used for proper offset calculation of intermediate data buffers + kernel.params.arguments.push_back({ArgumentDescriptor::Types::SCALAR, 0}); + } + const auto buf_sizes = get_internal_buffer_sizes(prim_params, kernel_idx); if (!prim_params.conf.is_paged_attention) { kd.internalBufferSizes.clear(); @@ -379,6 +396,15 @@ void SDPAKernelOpt::GetUpdateDispatchDataFunc(KernelData& kd) const { kernel_data.kernels[0].params.workGroups.global = dispatch_data.gws; kernel_data.kernels[0].params.workGroups.local = dispatch_data.lws; kernel_data.kernels[0].skip_execution = false; + + if (prim_params.outputs.size() > 1) { + const auto max_seq_len = prim_params.conf.paged_attention_max_len; + const auto seq_len_partition_size = get_seq_len_partition_size(params, prim_params.conf.head_size, KernelsTypes::MULTI_TOKENS); + + kernel_data.kernels[0].params.scalars.resize(1); + kernel_data.kernels[0].params.scalars[0].t = ScalarDescriptor::Types::UINT32; + kernel_data.kernels[0].params.scalars[0].v.u32 = static_cast(Align(max_seq_len, seq_len_partition_size)); + } } else { const auto num_of_partitions = get_partitions_num(prim_params, KernelsTypes::SINGLE_TOKEN); const auto buf_sizes = get_internal_buffer_sizes(prim_params, KernelsTypes::SINGLE_TOKEN); diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_opt.h b/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_opt.h index 8d7279f5546112..a4d351498d7075 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_opt.h +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_opt.h @@ -17,6 +17,9 @@ class SDPAKernelOpt : public SDPAKernelBase { KernelsPriority GetKernelsPriority(const Params& params) const override; ParamsKey GetSupportedKey() const override; + static size_t get_sg_number_scale_factor(const Params& params, size_t head_size, size_t kernel_type); + static size_t get_seq_len_partition_size(const Params& params, size_t head_size, size_t kernel_type); + protected: bool Validate(const Params& p) const override; void GetUpdateDispatchDataFunc(KernelData& kd) const override; diff --git a/src/plugins/intel_gpu/src/plugin/ops/paged_attention.cpp b/src/plugins/intel_gpu/src/plugin/ops/paged_attention.cpp index 7425b096b6d324..d82d3a66fed7f7 100644 --- a/src/plugins/intel_gpu/src/plugin/ops/paged_attention.cpp +++ b/src/plugins/intel_gpu/src/plugin/ops/paged_attention.cpp @@ -61,10 +61,13 @@ static void CreatePagedAttentionExtensionOp(ProgramBuilder& p, const std::shared OPENVINO_ASSERT(alibi_const != nullptr); prim.has_alibi = ov::shape_size(alibi_const->get_output_shape(0)) > 0; + prim.num_outputs = 1; if (op->get_output_size() > 1) { const auto scores_output_idx = 1; const auto& users = op->get_output_target_inputs(scores_output_idx); - OPENVINO_ASSERT(users.size() == 0, "[GPU] PagedAttention implementation doesn't support scores output yet"); + if (users.size() > 0) { + prim.num_outputs++; // Add scores output + } } p.add_primitive(*op, prim); diff --git a/src/plugins/intel_gpu/tests/unit/test_cases/paged_attention_gpu_test.cpp b/src/plugins/intel_gpu/tests/unit/test_cases/paged_attention_gpu_test.cpp new file mode 100644 index 00000000000000..a32ef3325cd9bc --- /dev/null +++ b/src/plugins/intel_gpu/tests/unit/test_cases/paged_attention_gpu_test.cpp @@ -0,0 +1,687 @@ +// Copyright (C) 2024 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "test_utils.h" +#include "random_generator.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include + +using namespace cldnn; +using namespace ov::intel_gpu; +using namespace ::tests; + +/* +* PagedAttention inputs: +* [0]: query +* shape: [batch_size_in_tokens, num_heads * head_size], type: f16 +* [1]: key +* shape: [batch_size_in_tokens, num_kv_heads * head_size], type: f16 +* [2]: value  +* shape: [batch_size_in_tokens, num_kv_heads * head_size], type: f16 +* [3]: key_cache +* shape: [num_blocks, num_kv_heads, head_size, block_size], type: f16 +* [4]: value_cache +* shape: [num_blocks, num_kv_heads, block_size, head_size], type: f16 +* [5]: past_lens +* shape: [batch_size_in_sequences], type: i32 +* [6]: subsequence_begins +* shape: [batch_size_in_sequences + 1], type: i32 +* [7]: block_indices +* Shape: [num_blocks], type: i32 +* [8]: block_indices_begins +* Shape: [batch_size_in_sequences + 1], type: i32 +* [9]: scale, optional +* [10]: sliding_window, optional +* [11]: alibi_slopes, optional +* [12]: max_context_len +* shape: [], type: i32 +*/ + +struct SubsequenceDescriptor { + int num_tokens; + int past_len; +}; + +struct PagedAttentionManager { + int num_heads; + int head_size; + int block_size; + std::vector subsequence_descs; + + // per-subsequence QKV inputs + std::vector> query_data; // {[1, num_tokens, num_heads, head_size], ..} + std::vector> key_data; // {[1, past_len + num_tokens, num_heads, head_size], ..} + std::vector> value_data; // {[1, past_len + num_tokens, num_heads, head_size], ..} + + // common PA inputs + std::vector past_lens; + std::vector subsequence_begins; + std::vector block_indices; + std::vector block_indices_begins; + std::vector max_context_len; + + cldnn::engine& test_engine; + cldnn::stream& test_stream; + tests::random_generator& rg; + + PagedAttentionManager(tests::random_generator& rg, + cldnn::engine& engine, + cldnn::stream& stream, + const std::vector& subsequence_descs, + int num_heads, + int head_size, + int block_size) + : num_heads(num_heads) + , head_size(head_size) + , block_size(block_size) + , subsequence_descs(subsequence_descs) + , test_engine(engine) + , test_stream(stream) + , rg(rg) { + // init subsequence_begins and block_indices_begins + subsequence_begins.push_back(0); + block_indices_begins.push_back(0); + + int max_len = 0; + for (int i = 0; i < static_cast(subsequence_descs.size()); i++) { + const auto& subsequence_desc = subsequence_descs[i]; + max_len = std::max(max_len, subsequence_desc.num_tokens + subsequence_desc.past_len); + + query_data.push_back(generate_input_data(rg, num_heads, subsequence_desc.num_tokens, head_size)); + key_data.push_back(generate_input_data(rg, num_heads, subsequence_desc.num_tokens + subsequence_desc.past_len, head_size)); + value_data.push_back(generate_input_data(rg, num_heads, subsequence_desc.num_tokens + subsequence_desc.past_len, head_size)); + + past_lens.push_back(subsequence_desc.past_len); + int subsequence_start_pos = subsequence_begins[i]; + int subsequence_end_pos = subsequence_start_pos + subsequence_desc.num_tokens; + subsequence_begins.push_back(subsequence_end_pos); + + int subsequence_length = subsequence_desc.num_tokens + subsequence_desc.past_len; + int required_blocks = ceil_div(subsequence_length, block_size); + int start_block_idx = block_indices.empty() ? 0 : block_indices.back() + 1; + int end_block_idx = start_block_idx + required_blocks; + for (int block_idx = start_block_idx; block_idx < end_block_idx; block_idx++) { + block_indices.push_back(block_idx); + } + + int block_indices_start_pos = block_indices_begins[i]; + int block_indices_end_pos = block_indices_start_pos + required_blocks; + block_indices_begins.push_back(block_indices_end_pos); + } + max_context_len.push_back(max_len); + } + + memory::ptr get_query_memory() { + return get_QKV_memory(query_data, false); + } + + memory::ptr get_key_memory() { + return get_QKV_memory(key_data, true); + } + + memory::ptr get_value_memory() { + return get_QKV_memory(value_data, true); + } + + memory::ptr get_key_cache_memory() { + auto num_blocks = block_indices.back() + 1; + auto key_cache_shape = ov::PartialShape{ num_blocks, num_heads, head_size, block_size }; + auto key_cache_layout = layout{ key_cache_shape, data_types::f16, format::bfyx }; + auto memory = test_engine.allocate_memory(key_cache_layout); + + for (int i = 0; i < static_cast(subsequence_descs.size()); i++) { + int past_len = subsequence_descs[i].past_len; + if (past_len != 0) { + int blocks_num = ceil_div(past_len, block_size); + int start_block_idx = block_indices[block_indices_begins[i]]; + for (int block_idx = 0; block_idx < blocks_num; block_idx++) { + int last_token_idx = block_idx == blocks_num - 1 ? past_len % block_size + : block_size; + for (int token_idx = 0; token_idx < last_token_idx; token_idx++) { + for (int head_idx = 0; head_idx < num_heads; head_idx++) { + for (int head_size_idx = 0; head_size_idx < head_size; head_size_idx++) { + size_t input_token_offset = block_idx * block_size + token_idx; + ov::float16* data_ptr = key_data[i].data() + + input_token_offset * num_heads * head_size + + head_idx * head_size + head_size_idx; + + // shape: [num_blocks, num_heads, head_size, block_size] + size_t output_offset = (start_block_idx + block_idx) * num_heads * head_size * block_size + + head_idx * head_size * block_size + + head_size_idx * block_size + + token_idx; + + set_values(test_stream, memory, data_ptr, 1, output_offset); + } + } + } + } + } + } + + return memory; + } + + memory::ptr get_value_cache_memory() { + auto num_blocks = block_indices.back() + 1; + auto value_cache_shape = ov::PartialShape{ num_blocks, num_heads, block_size, head_size }; + auto value_cache_layout = layout{ value_cache_shape, data_types::f16, format::bfyx }; + auto memory = test_engine.allocate_memory(value_cache_layout); + + for (int i = 0; i < static_cast(subsequence_descs.size()); i++) { + int past_len = subsequence_descs[i].past_len; + if (past_len != 0) { + int blocks_num = ceil_div(past_len, block_size); + int start_block_idx = block_indices[block_indices_begins[i]]; + for (int block_idx = 0; block_idx < blocks_num; block_idx++) { + int last_token_idx = block_idx == blocks_num - 1 ? past_len % block_size + : block_size; + for (int token_idx = 0; token_idx < last_token_idx; token_idx++) { + for (int head_idx = 0; head_idx < num_heads; head_idx++) { + size_t input_token_offset = block_idx * block_size + token_idx; + ov::float16* data_ptr = value_data[i].data() + + input_token_offset * num_heads * head_size + + head_idx * head_size; + + // shape: [num_blocks, num_heads, block_size, head_size] + size_t output_offset = (start_block_idx + block_idx) * num_heads * block_size * head_size + + head_idx * block_size * head_size + + token_idx * head_size; + + set_values(test_stream, memory, data_ptr, head_size, output_offset); + } + } + } + } + } + + return memory; + } + + memory::ptr get_past_lens_memory() { + return get_memory_from_vec(past_lens); + } + + memory::ptr get_subsequence_begins_memory() { + return get_memory_from_vec(subsequence_begins); + } + + memory::ptr get_block_indices_memory() { + return get_memory_from_vec(block_indices); + } + + memory::ptr get_block_indices_begins_memory() { + return get_memory_from_vec(block_indices_begins); + } + + memory::ptr get_scale_memory() { + std::vector scale = { ov::float16(get_default_scale()) }; + return get_memory_from_vec(scale); + } + + memory::ptr get_sliding_window_memory() { + std::vector sliding_window = { 0 }; + return get_memory_from_vec(sliding_window); + } + + memory::ptr get_alibi_memory() { + std::vector alibi; + return get_memory_from_vec(alibi); + } + + memory::ptr get_max_context_len_memory() { + return get_memory_from_vec(max_context_len); + } + + float get_default_scale() { + return static_cast(1.f / std::sqrt(head_size)); + } + +private: + template + memory::ptr get_memory_from_vec(std::vector& input_data) { + auto data_size = input_data.empty() ? 1 : input_data.size(); + auto shape = ov::PartialShape{ static_cast(data_size) }; + auto layout = cldnn::layout{ shape, ov::element::from(), format::bfyx }; + auto memory = test_engine.allocate_memory(layout); + + if (input_data.empty()) { + auto shape = ov::PartialShape{0}; + auto layout = cldnn::layout{ shape, ov::element::from(), format::bfyx }; + return test_engine.reinterpret_buffer(*memory, layout); + } + + set_values(test_stream, memory, input_data.data(), input_data.size(), 0); + + return memory; + } + + memory::ptr get_QKV_memory(std::vector>& input_data, bool skip_past_len) { + int total_tokens = 0; + for (const auto& subsequence_desc : subsequence_descs) + total_tokens += subsequence_desc.num_tokens; + + auto query_shape = ov::PartialShape{ total_tokens, num_heads * head_size }; + auto query_layout = layout{ query_shape, data_types::f16, format::bfyx }; + auto memory = test_engine.allocate_memory(query_layout); + + for (int subsequence_idx = 0; subsequence_idx < static_cast(subsequence_descs.size()); subsequence_idx++) { + for (int token_idx = 0; token_idx < subsequence_descs[subsequence_idx].num_tokens; token_idx++) { + for (int head_idx = 0; head_idx < num_heads; head_idx++) { + size_t input_token_offset = token_idx; + // as generated data stored in vectors includes past_len, ignore it for KV inputs + if (skip_past_len) + input_token_offset += subsequence_descs[subsequence_idx].past_len; + + ov::float16* data_ptr = input_data[subsequence_idx].data() + + input_token_offset * num_heads * head_size + + head_idx * head_size; + + size_t output_token_offset = subsequence_begins[subsequence_idx] + token_idx; + size_t output_offset = output_token_offset * num_heads * head_size + + head_idx * head_size; + + set_values(test_stream, memory, data_ptr, head_size, output_offset); + } + } + } + + return memory; + } + + template + static void set_values(stream& stream, memory::ptr mem, T* vals, size_t size, size_t dst_offset) { + mem_lock mem_ptr(mem, stream); + for (size_t i = 0; i < size; i++) { + mem_ptr[dst_offset + i] = vals[i]; + } + } + + static std::vector generate_input_data(tests::random_generator& rg, size_t num_heads, size_t tokens_num, size_t head_size) { + const size_t total_elements_num = tokens_num * num_heads * head_size; + auto data = rg.generate_random_1d(total_elements_num, -1, 1); + + return data; + } +}; + +struct PagedAttentionReference { + PagedAttentionReference(PagedAttentionManager& pam) + : pam(pam) + , test_engine(pam.test_engine) + , test_stream(pam.test_stream) {} + + std::pair, std::vector> get_reference() { + std::vector ref_data_output; + std::vector ref_scores_output; + + for (size_t i = 0; i < pam.subsequence_descs.size(); i++) { + const auto& subsequence_desc = pam.subsequence_descs[i]; + const auto kv_seq_len = subsequence_desc.num_tokens + subsequence_desc.past_len; + auto subsequence_ref_results = run_reference(pam.query_data[i], + pam.key_data[i], + pam.value_data[i], + subsequence_desc.num_tokens, + kv_seq_len, + pam.num_heads, + pam.head_size, + pam.get_default_scale()); + + // concatenate all subsequences into one vector + ref_data_output.insert(ref_data_output.end(), + subsequence_ref_results.first.begin(), + subsequence_ref_results.first.end()); + ref_scores_output.insert(ref_scores_output.end(), + subsequence_ref_results.second.begin(), + subsequence_ref_results.second.end()); + } + + return { ref_data_output, ref_scores_output }; + } + +private: + std::pair, std::vector> + run_reference(const std::vector& query_data, + const std::vector& key_data, + const std::vector& value_data, + int num_queries, + int num_keys, + int num_heads, + int head_size, + float scale) { + auto query_shape = ov::PartialShape{1, num_queries, num_heads, head_size}; + auto key_shape = ov::PartialShape{1, num_keys, num_heads, head_size}; + auto value_shape = ov::PartialShape{1, num_keys, num_heads, head_size}; + + auto query_layout = layout{query_shape, data_types::f16, format::bfyx}; + auto key_layout = layout{key_shape, data_types::f16, format::bfyx}; + auto value_layout = layout{value_shape, data_types::f16, format::bfyx}; + + OPENVINO_ASSERT(query_layout.count() == query_data.size()); + OPENVINO_ASSERT(key_layout.count() == key_data.size()); + OPENVINO_ASSERT(value_layout.count() == value_data.size()); + + auto query_mem = test_engine.allocate_memory(query_layout); + auto key_mem = test_engine.allocate_memory(key_layout); + auto value_mem = test_engine.allocate_memory(value_layout); + auto mask_mem = get_mask_mem(num_queries, num_keys, num_heads); + + set_values(query_mem, query_data); + set_values(key_mem, key_data); + set_values(value_mem, value_data); + + topology topology; + topology.add(input_layout("query", query_layout), + input_layout("key", key_layout), + input_layout("value", value_layout), + data("mask", mask_mem), + permute("query_transposed", input_info("query"), {0, 2, 1, 3}), + permute("key_transposed", input_info("key"), {0, 2, 1, 3}), + permute("value_transposed", input_info("value"), {0, 2, 1, 3}), + gemm("qk_gemm", { input_info("query_transposed"), input_info("key_transposed") }, data_types::f16, false, true, scale), + eltwise("eltwise", { input_info("qk_gemm"), input_info("mask") }, eltwise_mode::sum), + softmax("softmax", input_info("eltwise"), -1), + gemm("qkv_gemm", { input_info("softmax"), input_info("value_transposed") }, data_types::f16, false, false), + permute("qkv_gemm_transposed", input_info("qkv_gemm"), {0, 2, 1, 3}), + reorder("output_data", input_info("qkv_gemm_transposed"), format::bfyx, data_types::f16), + reorder("scores_data", input_info("softmax"), format::bfyx, data_types::f16) + ); + + ExecutionConfig config = get_test_default_config(test_engine); + config.set_property(ov::intel_gpu::optimize_data(true)); + config.set_property(ov::intel_gpu::allow_new_shape_infer(true)); + + network::ptr network = get_network(test_engine, topology, config, get_test_stream_ptr(), false); + network->set_input_data("query", query_mem); + network->set_input_data("key", key_mem); + network->set_input_data("value", value_mem); + + auto outputs = network->execute(); + + auto output_data_mem = outputs.at("output_data").get_memory(); + auto output_scores_mem = outputs.at("scores_data").get_memory(); + + return { get_output_data_vec(output_data_mem, num_queries, head_size, num_heads), + get_output_scores_vec(output_scores_mem, num_queries, num_keys, num_heads) }; + } + + std::vector get_output_scores_vec(memory::ptr scores_output, + int num_queries, + int num_keys, + int num_heads) { + OPENVINO_ASSERT(scores_output->count() == static_cast(num_heads * num_queries * num_keys)); + + std::vector output_scores(num_keys, 0); + mem_lock mem_ptr(scores_output, test_stream); + for (int head_idx = 0; head_idx < num_heads; head_idx++) { + for (int score_idx = 0; score_idx < num_keys; score_idx++) { + output_scores[score_idx] += mem_ptr[head_idx * num_queries * num_keys + + (num_queries - 1) * num_keys + + score_idx]; + } + } + + return output_scores; + } + + std::vector get_output_data_vec(memory::ptr data_output, + int num_queries, + int head_size, + int num_heads) { + OPENVINO_ASSERT(data_output->count() == static_cast(num_queries * num_heads * head_size)); + + std::vector output_data(data_output->count()); + mem_lock mem_ptr(data_output, test_stream); + for (size_t i = 0; i < data_output->count(); i++) + output_data[i] = mem_ptr[i]; + + return output_data; + } + + memory::ptr get_mask_mem(int num_queries, int num_keys, int num_heads) { + /* + * Two kinds of masks: + * + * Case 1 (N == K): + * num_queries = N + * num_keys = K = N + * head_size = H + * Q [N, H] * K[H, N] + * QK [N, N] + * 0 1 N + * 0 [ 0, MIN, .., MIN ] + * 1 [ 0, 0, .., MIN ] + * [ .., .., .., MIN ] + * N [ 0, 0, .., 0 ] + * + * Case 2 (N != K): + * num_queries = N + * num_keys = K + * head_size = H + * past_len = P = K - N + 1 + * Q [N, H] * K[H, K] + * QK [N, K] + * 0 1 2 P .. K + * 0 [ 0, 0, 0, MIN, MIN, MIN ] + * 1 [ 0, 0, 0, 0, MIN, MIN ] + * [ .., .., .., .., .., MIN ] + * N [ 0, 0, 0, 0, .., 0 ] + * + * Shapes: + * Q [1, num_heads, num_queries, head_size] + * K [1, num_heads, head_size, num_keys] + * Q*K [1, num_heads, num_queries, num_keys] + */ + + auto mask_shape = ov::PartialShape{ 1, 1, num_queries, num_keys }; + auto mask_layout = layout{mask_shape, data_types::f16, format::bfyx}; + auto mask_mem = test_engine.allocate_memory(mask_layout); + + int past_len = num_keys - num_queries + 1; + mem_lock mem_ptr(mask_mem, test_stream); + for (int i = 0; i < num_queries; i++) { + for (int j = 0; j < num_keys; j++) { + mem_ptr[i * num_keys + j] = j >= past_len + i ? std::numeric_limits::lowest() + : ov::float16(0.f); + } + } + + return mask_mem; + } + + + PagedAttentionManager& pam; + cldnn::engine& test_engine; + cldnn::stream& test_stream; +}; + +template +struct PagedAttentionTest : public ::testing::TestWithParam { +public: + random_generator rg; + cldnn::engine& engine = get_test_engine(); + float tolerance = 2e-3; + + void SetUp() override { + rg.set_seed(GET_SUITE_NAME); + } + + void execute(T& p) { + PagedAttentionManager pam(rg, get_test_engine(), get_test_stream(), p.subsequences, p.num_heads, p.head_size, p.block_size); + + auto query_mem = pam.get_query_memory(); + auto key_mem = pam.get_key_memory(); + auto value_mem = pam.get_value_memory(); + + auto key_cache_mem = pam.get_key_cache_memory(); + auto value_cache_mem = pam.get_value_cache_memory(); + + auto past_lens_mem = pam.get_past_lens_memory(); + auto subsequence_begins_mem = pam.get_subsequence_begins_memory(); + auto block_indices_mem = pam.get_block_indices_memory(); + auto block_indices_begins_mem = pam.get_block_indices_begins_memory(); + + auto scale_mem = pam.get_scale_memory(); + auto sliding_window_mem = pam.get_sliding_window_memory(); + auto alibi_mem = pam.get_alibi_memory(); + auto max_context_len_mem = pam.get_max_context_len_memory(); + + auto query_layout = query_mem->get_layout(); + auto key_layout = key_mem->get_layout(); + auto value_layout = value_mem->get_layout(); + auto key_cache_layout = key_cache_mem->get_layout(); + auto value_cache_layout = value_cache_mem->get_layout(); + auto past_lens_layout = past_lens_mem->get_layout(); + auto subsequence_begins_layout = subsequence_begins_mem->get_layout(); + auto block_indices_layout = block_indices_mem->get_layout(); + auto block_indices_begins_layout = block_indices_begins_mem->get_layout(); + auto scale_layout = scale_mem->get_layout(); + auto sliding_window_layout = sliding_window_mem->get_layout(); + auto alibi_layout = alibi_mem->get_layout(); + auto max_context_len_layout = max_context_len_mem->get_layout(); + + // make layouts dynamic + query_layout.set_partial_shape(ov::PartialShape{ -1, p.num_heads * p.head_size }); + key_layout.set_partial_shape(ov::PartialShape{ -1, p.num_heads * p.head_size }); + value_layout.set_partial_shape(ov::PartialShape{ -1, p.num_heads * p.head_size }); + key_cache_layout.set_partial_shape(ov::PartialShape{ -1, p.num_heads, p.head_size, p.block_size }); + value_cache_layout.set_partial_shape(ov::PartialShape{ -1, p.num_heads, p.block_size, p.head_size }); + past_lens_layout.set_partial_shape(ov::PartialShape{ -1 }); + subsequence_begins_layout.set_partial_shape(ov::PartialShape{ -1 }); + block_indices_layout.set_partial_shape(ov::PartialShape{ -1 }); + block_indices_begins_layout.set_partial_shape(ov::PartialShape{ -1 }); + + auto pa_prim = paged_attention("paged_attention", { input_info("query"), + input_info("key"), + input_info("value"), + input_info("key_cache"), + input_info("value_cache"), + input_info("past_lens"), + input_info("subsequence_begins"), + input_info("block_indices"), + input_info("block_indices_begins"), + input_info("scale"), + input_info("sliding_window"), + input_info("alibi"), + input_info("max_context_len") }); + + pa_prim.head_size = p.head_size; + pa_prim.kv_heads_num = p.num_heads; + pa_prim.heads_num = p.num_heads; + pa_prim.scale_val = pam.get_default_scale(); + pa_prim.has_alibi = false; + pa_prim.num_outputs = p.scores_output ? 2 : 1; + + topology topology; + topology.add( + input_layout("query", query_layout), + input_layout("key", key_layout), + input_layout("value", value_layout), + input_layout("key_cache", key_cache_layout), + input_layout("value_cache", value_cache_layout), + input_layout("past_lens", past_lens_layout), + input_layout("subsequence_begins", subsequence_begins_layout), + input_layout("block_indices", block_indices_layout), + input_layout("block_indices_begins", block_indices_begins_layout), + input_layout("scale", scale_layout), + input_layout("sliding_window", sliding_window_layout), + input_layout("alibi", alibi_layout), + input_layout("max_context_len", max_context_len_layout), + pa_prim, + reorder("output_data", input_info("paged_attention", 0), format::bfyx, data_types::f16) + ); + + if (p.scores_output) { + topology.add(reorder("output_scores", input_info("paged_attention", 1), format::bfyx, data_types::f16)); + } + + ExecutionConfig config = get_test_default_config(get_test_engine()); + config.set_property(ov::intel_gpu::optimize_data(true)); + config.set_property(ov::intel_gpu::allow_new_shape_infer(true)); + + network::ptr network = get_network(get_test_engine(), topology, config, get_test_stream_ptr(), false); + network->set_input_data("query", query_mem); + network->set_input_data("key", key_mem); + network->set_input_data("value", value_mem); + network->set_input_data("key_cache", key_cache_mem); + network->set_input_data("value_cache", value_cache_mem); + network->set_input_data("past_lens", past_lens_mem); + network->set_input_data("subsequence_begins", subsequence_begins_mem); + network->set_input_data("block_indices", block_indices_mem); + network->set_input_data("block_indices_begins", block_indices_begins_mem); + network->set_input_data("scale", scale_mem); + network->set_input_data("sliding_window", sliding_window_mem); + network->set_input_data("alibi", alibi_mem); + network->set_input_data("max_context_len", max_context_len_mem); + + auto outputs = network->execute(); + + cldnn::memory::ptr output_data_mem = nullptr; + cldnn::memory::ptr output_scores_mem = nullptr; + + output_data_mem = outputs.at("output_data").get_memory(); + if (p.scores_output) { + output_scores_mem = outputs.at("output_scores").get_memory(); + } + + auto ref_data = PagedAttentionReference(pam).get_reference(); + compare(output_data_mem, output_scores_mem, ref_data); + } + + void compare(memory::ptr data_output_mem, memory::ptr scores_output_mem, std::pair, std::vector> ref_data) { + if (data_output_mem) { + ASSERT_EQ(data_output_mem->count(), ref_data.first.size()); + mem_lock mem_ptr(data_output_mem, get_test_stream()); + for (size_t i = 0; i < data_output_mem->count(); i++) { + ASSERT_NEAR(mem_ptr[i], ref_data.first[i], tolerance); + } + } + + if (scores_output_mem) { + ASSERT_EQ(scores_output_mem->count(), ref_data.second.size()); + mem_lock mem_ptr(scores_output_mem, get_test_stream()); + for (size_t i = 0; i < scores_output_mem->count(); i++) { + ASSERT_NEAR(mem_ptr[i], ref_data.second[i], tolerance); + } + } + } +}; + +struct paged_attention_test_params { + std::vector subsequences; + int num_heads; + int head_size; + int block_size; + bool scores_output; +}; + +class paged_attention_test : public PagedAttentionTest {}; +TEST_P(paged_attention_test, basic) { + auto p = GetParam(); + + execute(p); +} + +INSTANTIATE_TEST_SUITE_P(smoke_paged_attention, paged_attention_test, ::testing::ValuesIn(std::vector{ + /* with scores output */ + paged_attention_test_params{ {{10, 0}}, 2, 64, 16, true }, // 1st token + paged_attention_test_params{ {{36, 0}}, 2, 64, 16, true }, // 1st token + paged_attention_test_params{ {{1024, 0}}, 2, 64, 16, true }, // 1st token long + paged_attention_test_params{ {{10, 0}, {30, 0}}, 2, 64, 16, true }, // 1st token + 1st token + paged_attention_test_params{ {{128, 0}, {256, 0}}, 2, 64, 16, true }, // 1st token + 1st token + paged_attention_test_params{ {{1, 10}}, 2, 64, 16, true }, // 2nd token + paged_attention_test_params{ {{1, 34}, {1, 515}}, 2, 64, 16, true }, // 2nd token + 2nd token + paged_attention_test_params{ {{1, 34}, {25, 0}, {10, 34}}, 2, 64, 16, true }, // mixed: 2nd token + 1st token + part of 1st token + /* without scores output */ + paged_attention_test_params{ {{10, 0}}, 2, 64, 16, false }, // 1st token + paged_attention_test_params{ {{1024, 0}}, 2, 64, 16, false }, // 1st token long + paged_attention_test_params{ {{1, 34}, {1, 515}}, 2, 64, 16, false }, // 2nd token + 2nd token +})); diff --git a/src/plugins/intel_npu/src/plugin/npuw/llm_compiled_model.cpp b/src/plugins/intel_npu/src/plugin/npuw/llm_compiled_model.cpp index 260a1c444284cb..eb13bc8b5bd1d9 100644 --- a/src/plugins/intel_npu/src/plugin/npuw/llm_compiled_model.cpp +++ b/src/plugins/intel_npu/src/plugin/npuw/llm_compiled_model.cpp @@ -5,14 +5,208 @@ #include "llm_infer_request.hpp" #include "logging.hpp" +#include "openvino/op/ops.hpp" +#include "openvino/openvino.hpp" +#include "openvino/opsets/opset13.hpp" +#include "openvino/pass/graph_rewrite.hpp" +#include "openvino/pass/matcher_pass.hpp" +#include "openvino/pass/pattern/op/wrap_type.hpp" #include "openvino/pass/stateful_to_stateless.hpp" +#include "openvino/pass/validate.hpp" #include "openvino/runtime/iasync_infer_request.hpp" +namespace opp = ov::pass::pattern; +class TransposeValueTensors : public ov::pass::MatcherPass { +public: + struct Context { + std::vector> new_params; + std::vector> old_params; + using Ref = std::reference_wrapper; + }; + + OPENVINO_MATCHER_PASS_RTTI("npuw::LLMCompiledModel::TransposeValueTensors"); + TransposeValueTensors(Context::Ref ctx) { + auto param = opp::wrap_type(); + auto transpose = opp::wrap_type({opp::any_input(), opp::any_input()}); + auto concat = opp::wrap_type({param, transpose}); + auto softmax = opp::wrap_type({opp::any_input()}); + auto matmul = opp::wrap_type({softmax, concat}); + + auto callback = [=](ov::pass::pattern::Matcher& m) { + auto& node_to_output = m.get_pattern_value_map(); + + auto matched_node_param = node_to_output.at(param).get_node_shared_ptr(); + auto matched_node_concat = node_to_output.at(concat).get_node_shared_ptr(); + auto matched_node_transpose = node_to_output.at(transpose).get_node_shared_ptr(); + auto matched_node_matmul = node_to_output.at(matmul).get_node_shared_ptr(); + + auto matched_param = std::static_pointer_cast(matched_node_param); + auto matched_concat = std::static_pointer_cast(matched_node_concat); + auto matched_transpose = std::static_pointer_cast(matched_node_transpose); + auto matched_matmul = std::static_pointer_cast(matched_node_matmul); + + auto shape = matched_param->get_partial_shape(); + OPENVINO_ASSERT(shape.size() == 4u); + // NB: Transpose Parameter that correspond to V-tensor it will + // speed-up its multiplication with attention scores + std::swap(shape[2], shape[3]); + auto new_param = std::make_shared(matched_param->get_element_type(), shape); + new_param->set_friendly_name(matched_param->get_friendly_name()); + new_param->outputs().begin()->get_tensor().set_names( + matched_param->outputs().begin()->get_tensor().get_names()); + ov::replace_node(matched_param, new_param); + // NB: Save in order to add/remove to the model later on + ctx.get().new_params.push_back(new_param); + ctx.get().old_params.push_back(matched_param); + + auto order_cst = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{4}, {0, 2, 3, 1}); + auto new_transpose = + std::make_shared(matched_transpose->input_value(0), order_cst->output(0)); + new_transpose->set_friendly_name(matched_transpose->get_friendly_name()); + ov::replace_node(matched_transpose, new_transpose); + + auto new_concat = + std::make_shared(ov::OutputVector{new_param->output(0), new_transpose->output(0)}, + 3u); + new_concat->set_friendly_name(matched_concat->get_friendly_name()); + ov::replace_node(matched_concat, new_concat); + + matched_matmul->set_transpose_b(true); + + return true; + }; + register_matcher(std::make_shared(matmul, "TransposeValueTensors"), std::move(callback)); + } +}; + +class ScaledDotProductAttentionDecomposition : public ov::pass::MatcherPass { +public: + OPENVINO_MATCHER_PASS_RTTI("npuw::LLMCompiledModel::ScaledDotProductAttentionDecomposition"); + ScaledDotProductAttentionDecomposition() { + auto pattern_node = ov::pass::pattern::wrap_type(); + + ov::matcher_pass_callback callback = [=](ov::pass::pattern::Matcher& m) { + auto& pattern_to_output = m.get_pattern_value_map(); + auto node = ov::as_type_ptr( + pattern_to_output.at(pattern_node).get_node_shared_ptr()); + + if (node == nullptr || transformation_callback(node)) { + return false; + } + + auto new_output_node = decompose(node); + ov::replace_node(node, new_output_node); + return true; + }; + + auto m = std::make_shared(pattern_node, "ScaledDotProductAttentionDecomposition"); + register_matcher(m, std::move(callback)); + } + std::shared_ptr decompose(std::shared_ptr node) { + using namespace ov::op; + using namespace ov; + auto query = node->input_value(0); + auto key = node->input_value(1); + auto value = node->input_value(2); + auto q_shape = register_new_node(query, element::i32); + auto k_shape = register_new_node(key, element::i32); + auto minus_one = register_new_node(v0::Constant::create(element::i32, Shape{}, {-1})); + auto minus_two = register_new_node(v0::Constant::create(element::i32, Shape{}, {-2})); + auto zero_i = register_new_node(v0::Constant::create(element::i32, Shape{}, {0})); + auto one_i = register_new_node(v0::Constant::create(element::i32, Shape{}, {1})); + auto one_f = register_new_node(one_i, query); + auto zero_f = register_new_node(zero_i, query); + + Output scale; + if (node->get_input_size() < 5) { + scale = register_new_node(q_shape, minus_one, zero_i)->output(0); + scale = register_new_node(scale, query); + auto sqrt_scale = register_new_node(scale); + scale = register_new_node(one_f, sqrt_scale); + } else { + scale = node->input_value(4); + } + + auto q_scaled = register_new_node(query, scale); + auto k_rank = register_new_node(k_shape, element::i32)->output(0); + auto k_last_dim = register_new_node(k_rank, minus_one); + auto k_next_dim = register_new_node(k_rank, minus_two)->output(0); + k_rank = register_new_node(k_rank, zero_i); + auto minus_inf = + register_new_node(v0::Constant::create(element::f32, Shape{}, {-std::numeric_limits::infinity()})) + ->output(0); + auto keep_dim_last = register_new_node(k_next_dim, zero_i); + auto k_dims_before_transpose = register_new_node(zero_i, keep_dim_last, one_i, element::i32); + + auto scaled_atten = register_new_node(q_scaled, key, false, true)->output(0); + minus_inf = register_new_node(minus_inf, scaled_atten); + + if (node->get_causal() || node->get_input_size() > 3) { + Output mask; + Output atten_mask; + if (!node->get_causal()) { + mask = node->input_value(3); + + // two types of masks are supported. A boolean mask where a value of True indicates that the element + // should take part in attention. A float mask of the same type as query, key, value that is added to + // the attention score. + if (mask.get_element_type() == element::boolean) { + atten_mask = register_new_node(mask, scaled_atten); + auto inv_mask = register_new_node(mask); + atten_mask = register_new_node(inv_mask, atten_mask, minus_inf); + } else { + atten_mask = mask; + } + } else { + auto target_s_len = register_new_node(q_shape, minus_two, zero_i); + auto source_s_len = register_new_node(k_shape, minus_two, zero_i); + auto ssl = register_new_node(source_s_len, zero_i); + auto tsl = register_new_node(target_s_len, zero_i); + auto mask_shape = register_new_node(OutputVector{tsl, ssl}, 0); + mask = register_new_node(minus_inf, mask_shape); + auto horizontal_range = + register_new_node(zero_i, source_s_len, one_i, element::i32)->output(0); + horizontal_range = register_new_node(horizontal_range, zero_i); + auto stop = register_new_node(target_s_len, one_i); + auto vertical_range = register_new_node(one_i, stop, one_i, element::i32)->output(0); + vertical_range = register_new_node(vertical_range, one_i); + auto triu = register_new_node(horizontal_range, vertical_range); + atten_mask = register_new_node(triu, mask, zero_f); + } + scaled_atten = register_new_node(scaled_atten, atten_mask); + } + + scaled_atten = register_new_node(scaled_atten, -1); + auto result = register_new_node(scaled_atten, value); + result->set_friendly_name(node->get_friendly_name()); + copy_runtime_info(node, get_new_nodes()); + return result; + } +}; + namespace { uint32_t align_to(uint32_t value, uint32_t alignment) { return (value + alignment - 1) & ~(alignment - 1); } +std::shared_ptr cvt_kvcache_to_fp16(const std::shared_ptr& model) { + ov::preprocess::PrePostProcessor ppp(model); + + for (const auto& tensor : model->inputs()) { + if (tensor.get_any_name().find("past_key") != std::string::npos) { + ppp.input(tensor.get_any_name()).tensor().set_element_type(ov::element::Type_t::f16); + } + } + + for (const auto& tensor : model->outputs()) { + if (tensor.get_any_name().find("present") != std::string::npos) { + ppp.output(tensor.get_any_name()).tensor().set_element_type(ov::element::Type_t::f16); + } + } + + return ppp.build(); +} + std::shared_ptr redirect_new_kv_to_output(const std::shared_ptr& model) { const auto kStartOutputKVCacheLayers = 1u; for (std::size_t i = kStartOutputKVCacheLayers; i < model->outputs().size(); ++i) { @@ -27,22 +221,33 @@ std::shared_ptr redirect_new_kv_to_output(const std::shared_ptr cvt_kvcache_to_fp16(const std::shared_ptr& model) { +std::shared_ptr cvt_value_tensors_layout(std::shared_ptr model) { ov::preprocess::PrePostProcessor ppp(model); - - for (const auto& tensor : model->inputs()) { - if (tensor.get_any_name().find("past_key") != std::string::npos) { - ppp.input(tensor.get_any_name()).tensor().set_element_type(ov::element::Type_t::f16); + for (auto tensor : model->outputs()) { + if (tensor.get_any_name().find("value") != std::string::npos) { + // NB: [batch, num_heads, seq_len, emb_size] -> [batch, num_heads, emb_size, seq_len] + ppp.output(tensor.get_any_name()).model().set_layout(ov::Layout("BHSE")); + ppp.output(tensor.get_any_name()).tensor().set_layout(ov::Layout("BHES")); } } + return ppp.build(); +} - for (const auto& tensor : model->outputs()) { - if (tensor.get_any_name().find("present") != std::string::npos) { - ppp.output(tensor.get_any_name()).tensor().set_element_type(ov::element::Type_t::f16); - } +bool optimize_value_tensors(std::shared_ptr model) { + ov::pass::GraphRewrite rewr; + rewr.add_matcher(); + TransposeValueTensors::Context ctx; + rewr.add_matcher(std::ref(ctx)); + rewr.run_on_model(model); + + model->add_parameters(ctx.new_params); + for (auto old_param : ctx.old_params) { + model->remove_parameter(old_param); } + ov::pass::Validate().run_on_model(model); - return ppp.build(); + // NB: if new_params is not empty - pass has been applied + return !ctx.new_params.empty(); } struct KVAxesPosition { @@ -116,32 +321,6 @@ std::optional extract_npu_descriptor(const std::shared_ptr(), max_tiles.as()}); } -std::optional pop_option(ov::AnyMap& config, const std::string& option_name) { - if (auto it = config.find(option_name); it != config.end()) { - std::optional found = std::make_optional(it->second); - config.erase(it); - return found; - } - return std::nullopt; -} - -template -std::optional get_option(ov::AnyMap& config, const std::string& option_name) { - if (auto it = config.find(option_name); it != config.end()) { - return std::make_optional(it->second.as()); - } - return std::nullopt; -} - -template -T pop_or_default(ov::AnyMap& config, const std::string& key, const T& default_value) { - auto anyopt = pop_option(config, key); - if (anyopt.has_value()) { - return anyopt.value().as(); - } - return default_value; -} - ov::AnyMap get_baseline_common_config() { ov::AnyMap config = { {"NPU_COMPILATION_MODE_PARAMS", "compute-layers-with-higher-precision=Sqrt,Power,ReduceMean,Add_RMSNorm"}, @@ -206,12 +385,6 @@ void merge_config_with(ov::AnyMap& lhs, const ov::AnyMap& rhs) { } } -void drop_cache_dir(ov::AnyMap& config) { - if (config.count("NPU_USE_NPUW") != 0u) { - pop_option(config, "CACHE_DIR"); - } -} - void split_llm_properties(const ov::AnyMap& properties, ov::AnyMap& llm_properties, ov::AnyMap& other_properties) { for (auto it = properties.begin(); it != properties.end(); ++it) { if (it->first.find("NPUW_LLM") != it->first.npos) { @@ -251,41 +424,48 @@ ov::npuw::LLMCompiledModel::LLMCompiledModel(const std::shared_ptr& m auto kvcache_model = model->clone(); LOG_DEBUG("2. Transform kvcache model from stateful to stateless."); ov::pass::StatefulToStateless().run_on_model(kvcache_model); - LOG_DEBUG("3. Creating prefill model as clone of transformed kvcache one."); auto prefill_model = kvcache_model->clone(); prefill_model->set_friendly_name(kvcache_model->get_friendly_name() + "_prefill"); - LOG_DEBUG("4. Converting KV-cache in prefill model to FP16."); - prefill_model = cvt_kvcache_to_fp16(prefill_model); - - LOG_DEBUG("5. Optimize kvcache kvcache model to output key/values for new token."); - kvcache_model = redirect_new_kv_to_output(kvcache_model); - LOG_DEBUG("6. Converting KV-cache in kvcache model to FP16."); - kvcache_model = cvt_kvcache_to_fp16(kvcache_model); + const ::intel_npu::npuw::llm::ModelDesc model_desc = m_cfg.get<::intel_npu::NPUW_LLM_MODEL_DESC>(); const uint32_t kMaxPromptLen = align_to(m_cfg.get<::intel_npu::NPUW_LLM_MAX_PROMPT_LEN>(), 64u); const uint32_t kMinResponseLen = align_to(m_cfg.get<::intel_npu::NPUW_LLM_MIN_RESPONSE_LEN>(), 64u); - const ::intel_npu::npuw::llm::ModelDesc model_desc = m_cfg.get<::intel_npu::NPUW_LLM_MODEL_DESC>(); KVAxesPosition axes = get_kv_axes(model_desc.type); m_kvcache_desc = KVCacheDesc{kMaxPromptLen, kMaxPromptLen + kMinResponseLen, 0u, axes.seq_len}; - LOG_DEBUG("7. Make prefill model with static shapes"); + LOG_DEBUG("4. Make prefill model with static shapes"); reshape_to_static(prefill_model, m_kvcache_desc.max_prompt_size, m_kvcache_desc.max_prompt_size, axes); - LOG_DEBUG("8. Make kvcache model with static shapes"); + LOG_DEBUG("5. Make kvcache model with static shapes"); reshape_to_static(kvcache_model, 1u, m_kvcache_desc.total_size, axes); + LOG_DEBUG("6.Check and apply opt layout if applicable."); + // NB: Try to apply opt transpose only for Llama-2-7b-chat-hf model + if (model_desc.name_or_path == "meta-llama/Llama-2-7b-chat-hf" || + (model_desc.type == "llama" && model_desc.num_key_value_heads == 32)) { + if (optimize_value_tensors(kvcache_model)) { + // NB: Check if TransposeValueTensors transformation was applied + m_kvcache_desc.v_tensors_transposed = true; + prefill_model = cvt_value_tensors_layout(prefill_model); + } + } + LOG_DEBUG("7. Optimize kvcache model to output key/values for new token."); + kvcache_model = redirect_new_kv_to_output(kvcache_model); + LOG_DEBUG("8. Converting KV-cache in kvcache model to FP16."); + kvcache_model = cvt_kvcache_to_fp16(kvcache_model); + LOG_DEBUG("9. Converting KV-cache in prefill model to FP16."); + prefill_model = cvt_kvcache_to_fp16(prefill_model); auto npudesc = extract_npu_descriptor(plugin); - - ov::AnyMap properties_copy = std::move(other_props); + ov::AnyMap properties_copy = other_props; auto prefill_config = get_default_prefill_config(model, npudesc); + // NB: GENERATE_HINT is only applicable for default generate config! const ::intel_npu::npuw::llm::GenerateHint generate_hint = m_cfg.get<::intel_npu::NPUW_LLM_GENERATE_HINT>(); - LOG_DEBUG("9. Passed GENERATE_HINT: " << std::string(::intel_npu::NPUW_LLM_GENERATE_HINT::toString(generate_hint))); + LOG_DEBUG( + "10. Passed GENERATE_HINT: " << std::string(::intel_npu::NPUW_LLM_GENERATE_HINT::toString(generate_hint))); auto generate_config = get_default_generate_config(model, npudesc, generate_hint); + merge_config_with(prefill_config, properties_copy); merge_config_with(generate_config, properties_copy); - // FIXME: Drop CACHE_DIR option if NPUW is enabled - drop_cache_dir(prefill_config); - drop_cache_dir(generate_config); m_kvcache_compiled = std::make_shared(kvcache_model, plugin, generate_config); m_prefill_compiled = std::make_shared(prefill_model, plugin, prefill_config); diff --git a/src/plugins/intel_npu/src/plugin/npuw/llm_compiled_model.hpp b/src/plugins/intel_npu/src/plugin/npuw/llm_compiled_model.hpp index 1a748997fd48fa..e37a47b2c77948 100644 --- a/src/plugins/intel_npu/src/plugin/npuw/llm_compiled_model.hpp +++ b/src/plugins/intel_npu/src/plugin/npuw/llm_compiled_model.hpp @@ -22,6 +22,7 @@ class LLMCompiledModel : public ov::npuw::ICompiledModel { uint32_t total_size = 0u; uint32_t num_stored_tokens = 0u; uint32_t dim = 0u; + bool v_tensors_transposed = false; }; LLMCompiledModel(const std::shared_ptr& model, diff --git a/src/plugins/intel_npu/src/plugin/npuw/llm_infer_request.cpp b/src/plugins/intel_npu/src/plugin/npuw/llm_infer_request.cpp index a73478c0cab5d2..12f103cc0ab6a2 100644 --- a/src/plugins/intel_npu/src/plugin/npuw/llm_infer_request.cpp +++ b/src/plugins/intel_npu/src/plugin/npuw/llm_infer_request.cpp @@ -27,6 +27,36 @@ ov::SoPtr make_tensor_slice(ov::SoPtr tensor, end_shape[dim] = end_pos; return ov::get_tensor_impl(ov::Tensor(ov::make_tensor(tensor), start_shape, end_shape)); } + +void copy_columns_by_row_chunks(ov::SoPtr src, ov::SoPtr& dst) { + const auto src_shape = src->get_shape(); + + OPENVINO_ASSERT(src_shape.size() == 4u); + OPENVINO_ASSERT(src_shape == dst->get_shape()); + OPENVINO_ASSERT(src->get_byte_size() == dst->get_byte_size()); + + const auto src_strides = src->get_strides(); + const auto dst_strides = dst->get_strides(); + const auto elem_size = src->get_byte_size() / src->get_size(); + + const auto C = src_shape[1]; + const auto H = src_shape[2]; + const auto W = src_shape[3]; + + const auto IS_H = src_strides[2]; + const auto OS_H = dst_strides[2]; + + const size_t chunk_byte_size = W * elem_size; + + const auto* src_p = static_cast(src->data()); + auto* dst_p = static_cast(dst->data()); + + for (size_t i = 0; i < C * H; ++i) { + const size_t src_offset = i * IS_H; + const size_t dst_offset = i * OS_H; + std::copy_n(src_p + src_offset, chunk_byte_size, dst_p + dst_offset); + } +} } // anonymous namespace ov::npuw::LLMInferRequest::LLMInferRequest(const std::shared_ptr& compiled_model, @@ -116,17 +146,25 @@ void ov::npuw::LLMInferRequest::infer_generate(ov::SoPtr input_ids, // taking into account kvcache dimension. fill_tensor(kvcache_in_tensor, 0); + const auto& kv_dim = (output_name.find("value") != std::string::npos && m_kvcache_desc.v_tensors_transposed) + ? 3u + : m_kvcache_desc.dim; + auto prefill_out_slice = make_tensor_slice(prefill_out_tensor, - m_kvcache_desc.dim, + kv_dim, m_kvcache_desc.max_prompt_size - m_kvcache_desc.num_stored_tokens, m_kvcache_desc.max_prompt_size); - auto kvcache_in_slice = - make_tensor_slice(kvcache_in_tensor, m_kvcache_desc.dim, 0u, m_kvcache_desc.num_stored_tokens); + auto kvcache_in_slice = make_tensor_slice(kvcache_in_tensor, kv_dim, 0u, m_kvcache_desc.num_stored_tokens); - prefill_out_slice->copy_to(kvcache_in_slice._ptr); + if (kv_dim == 3u) { + copy_columns_by_row_chunks(prefill_out_slice, kvcache_in_slice); + } else { + prefill_out_slice->copy_to(kvcache_in_slice._ptr); + } } + LOG_DEBUG("Prepare attention mask pattern."); auto* attention_mask_data = m_kvcache_request->get_tensor(m_kvcache_in_ports.at("attention_mask"))->data(); @@ -156,8 +194,11 @@ void ov::npuw::LLMInferRequest::infer_generate(ov::SoPtr input_ids, const auto& output_name = kvcache_compiled->outputs()[kStartOutputKVCacheLayers + i].get_any_name(); const auto& input_name = std::regex_replace(output_name, std::regex("present"), "past_key_values"); auto kvcache_in_tensor = m_kvcache_request->get_tensor(m_kvcache_in_ports.at(input_name)); + const auto& kv_dim = (output_name.find("value") != std::string::npos && m_kvcache_desc.v_tensors_transposed) + ? 3u + : m_kvcache_desc.dim; auto kvcache_in_slice = make_tensor_slice(kvcache_in_tensor, - m_kvcache_desc.dim, + kv_dim, m_kvcache_desc.num_stored_tokens - 1, m_kvcache_desc.num_stored_tokens); auto kvcache_out_tensor = m_kvcache_request->get_tensor(m_kvcache_out_ports.at(output_name)); diff --git a/src/tests/functional/plugin/shared/include/single_op_tests/fake_convert.hpp b/src/tests/functional/plugin/shared/include/single_op_tests/fake_convert.hpp new file mode 100644 index 00000000000000..d22809e332b0a3 --- /dev/null +++ b/src/tests/functional/plugin/shared/include/single_op_tests/fake_convert.hpp @@ -0,0 +1,16 @@ +// Copyright (C) 2024 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#pragma once + +#include "shared_test_classes/single_op/fake_convert.hpp" + +namespace ov { +namespace test { + +TEST_P(FakeConvertLayerTest, Inference) { + run(); +} +} // namespace test +} // namespace ov diff --git a/src/tests/functional/shared_test_classes/include/shared_test_classes/single_op/fake_convert.hpp b/src/tests/functional/shared_test_classes/include/shared_test_classes/single_op/fake_convert.hpp new file mode 100644 index 00000000000000..ce6ad97aba1b5d --- /dev/null +++ b/src/tests/functional/shared_test_classes/include/shared_test_classes/single_op/fake_convert.hpp @@ -0,0 +1,28 @@ +// Copyright (C) 2024 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#pragma once + +#include "shared_test_classes/base/ov_subgraph.hpp" + +namespace ov { +namespace test { +using FakeConvertParams = std::tuple, // Data shape + Shape, // Scale shape + Shape, // Shift shape + ov::element::Type, // Input precision + ov::element::Type, // Ddestination precision + bool, // Default shift + std::string>; // Device name + +class FakeConvertLayerTest : public testing::WithParamInterface, + virtual public ov::test::SubgraphBaseTest { +public: + static std::string getTestCaseName(const testing::TestParamInfo& obj); + +protected: + void SetUp() override; +}; +} // namespace test +} // namespace ov diff --git a/src/tests/functional/shared_test_classes/src/single_op/fake_convert.cpp b/src/tests/functional/shared_test_classes/src/single_op/fake_convert.cpp new file mode 100644 index 00000000000000..d207a8dabfb883 --- /dev/null +++ b/src/tests/functional/shared_test_classes/src/single_op/fake_convert.cpp @@ -0,0 +1,64 @@ +// Copyright (C) 2024 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "shared_test_classes/single_op/fake_convert.hpp" + +#include "openvino/opsets/opset1.hpp" +#include "openvino/opsets/opset13.hpp" + +namespace ov { +namespace test { +std::string FakeConvertLayerTest::getTestCaseName(const testing::TestParamInfo& obj) { + FakeConvertParams params = obj.param; + + std::vector data_shapes; + Shape scale_shape, shift_shape; + element::Type_t data_prec, dst_prec; + bool default_shift; + std::string target_device; + std::tie(data_shapes, scale_shape, shift_shape, data_prec, dst_prec, default_shift, target_device) = params; + + std::ostringstream result; + result << "IS=("; + for (const auto& shape : data_shapes) { + result << ov::test::utils::partialShape2str({shape.first}) << "_"; + } + result << ")_TS=("; + for (const auto& shape : data_shapes) { + for (const auto& item : shape.second) { + result << ov::test::utils::vec2str(item) << "_"; + } + } + result << ")_scaleShape=" << ov::test::utils::vec2str(scale_shape) << "_"; + result << "shiftShape=" << ov::test::utils::vec2str(shift_shape) << "_"; + result << "dataPrecision=" << element::Type(data_prec) << "_"; + result << "destinationPrecision=" << element::Type(dst_prec) << "_"; + if (default_shift) + result << "defaultShift=true"; + else + result << "defaultShift=false"; + return result.str(); +} + +void FakeConvertLayerTest::SetUp() { + FakeConvertParams params = this->GetParam(); + + std::vector data_shapes; + Shape scale_shape, shift_shape; + element::Type_t data_prec, dst_prec; + bool default_shift; + std::tie(data_shapes, scale_shape, shift_shape, data_prec, dst_prec, default_shift, targetDevice) = params; + + init_input_shapes(data_shapes); + + const auto data = std::make_shared(data_prec, inputDynamicShapes.front()); + const auto scale = std::make_shared(data_prec, scale_shape); + const auto shift = std::make_shared(data_prec, shift_shape); + + const auto fake_convert = default_shift ? std::make_shared(data, scale, dst_prec) + : std::make_shared(data, scale, shift, dst_prec); + function = std::make_shared(NodeVector{fake_convert}, ParameterVector{data}); +} +} // namespace test +} // namespace ov diff --git a/tests/constraints.txt b/tests/constraints.txt index 4f46cd0cc8b2e9..c339ac3c65d56f 100644 --- a/tests/constraints.txt +++ b/tests/constraints.txt @@ -21,11 +21,8 @@ pytest>=5.0,<8.4 pytest-dependency==0.5.1 pytest-html==4.1.1 pytest-timeout==2.3.1 -jax<=0.4.36 -jaxlib<=0.4.36 kornia==0.7.0 networkx<=3.3 -flax<=0.10.2 --extra-index-url https://download.pytorch.org/whl/cpu torch~=2.5.1; platform_system != "Darwin" or platform_machine != "x86_64" diff --git a/tests/layer_tests/onnx_tests/test_abs.py b/tests/layer_tests/onnx_tests/test_abs.py index 9a82929ea35547..71e509faef3e65 100644 --- a/tests/layer_tests/onnx_tests/test_abs.py +++ b/tests/layer_tests/onnx_tests/test_abs.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_and.py b/tests/layer_tests/onnx_tests/test_and.py index ca5d21a42fe067..195ace1dadfa14 100644 --- a/tests/layer_tests/onnx_tests/test_and.py +++ b/tests/layer_tests/onnx_tests/test_and.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_argmax.py b/tests/layer_tests/onnx_tests/test_argmax.py index 604df5e7e69875..80d7568e9e8c4c 100644 --- a/tests/layer_tests/onnx_tests/test_argmax.py +++ b/tests/layer_tests/onnx_tests/test_argmax.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_ceil.py b/tests/layer_tests/onnx_tests/test_ceil.py index b7558630ac1c63..ea7ea10abbd31d 100644 --- a/tests/layer_tests/onnx_tests/test_ceil.py +++ b/tests/layer_tests/onnx_tests/test_ceil.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_clip.py b/tests/layer_tests/onnx_tests/test_clip.py index dbce45193034d9..3cb3ba250a12e0 100644 --- a/tests/layer_tests/onnx_tests/test_clip.py +++ b/tests/layer_tests/onnx_tests/test_clip.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_concat.py b/tests/layer_tests/onnx_tests/test_concat.py index 8627f3b198dbd3..602b6a69644527 100644 --- a/tests/layer_tests/onnx_tests/test_concat.py +++ b/tests/layer_tests/onnx_tests/test_concat.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_conv.py b/tests/layer_tests/onnx_tests/test_conv.py index b7f9729141c33e..202d6af2915c67 100644 --- a/tests/layer_tests/onnx_tests/test_conv.py +++ b/tests/layer_tests/onnx_tests/test_conv.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_cumsum.py b/tests/layer_tests/onnx_tests/test_cumsum.py index 1e197de490d518..486b1f50835fb0 100644 --- a/tests/layer_tests/onnx_tests/test_cumsum.py +++ b/tests/layer_tests/onnx_tests/test_cumsum.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_dequantize_linear.py b/tests/layer_tests/onnx_tests/test_dequantize_linear.py index 9090f3a829919b..319030590a3f0d 100644 --- a/tests/layer_tests/onnx_tests/test_dequantize_linear.py +++ b/tests/layer_tests/onnx_tests/test_dequantize_linear.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_elu.py b/tests/layer_tests/onnx_tests/test_elu.py index dbffc32d09c6c7..9f0321ec9a6ee3 100644 --- a/tests/layer_tests/onnx_tests/test_elu.py +++ b/tests/layer_tests/onnx_tests/test_elu.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_embedding_bag.py b/tests/layer_tests/onnx_tests/test_embedding_bag.py index a18a59b9752f16..54d940c01fb36c 100644 --- a/tests/layer_tests/onnx_tests/test_embedding_bag.py +++ b/tests/layer_tests/onnx_tests/test_embedding_bag.py @@ -5,6 +5,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + import torch import torch.nn as nn from common.layer_test_class import CommonLayerTest, check_ir_version diff --git a/tests/layer_tests/onnx_tests/test_floor.py b/tests/layer_tests/onnx_tests/test_floor.py index 87ad058c510e8c..5076befc414941 100644 --- a/tests/layer_tests/onnx_tests/test_floor.py +++ b/tests/layer_tests/onnx_tests/test_floor.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_gather.py b/tests/layer_tests/onnx_tests/test_gather.py index a45d5b4f4a916b..9380de31c6dccc 100644 --- a/tests/layer_tests/onnx_tests/test_gather.py +++ b/tests/layer_tests/onnx_tests/test_gather.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_hard_sigmoid.py b/tests/layer_tests/onnx_tests/test_hard_sigmoid.py index 12986c590d41d4..a62ab2a7fc54e8 100644 --- a/tests/layer_tests/onnx_tests/test_hard_sigmoid.py +++ b/tests/layer_tests/onnx_tests/test_hard_sigmoid.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_identity.py b/tests/layer_tests/onnx_tests/test_identity.py index a86c0e2a687257..e58e272de49ec0 100644 --- a/tests/layer_tests/onnx_tests/test_identity.py +++ b/tests/layer_tests/onnx_tests/test_identity.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_leaky_relu.py b/tests/layer_tests/onnx_tests/test_leaky_relu.py index 3a12bfcd92c33e..cff9cd87b59d30 100644 --- a/tests/layer_tests/onnx_tests/test_leaky_relu.py +++ b/tests/layer_tests/onnx_tests/test_leaky_relu.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_log.py b/tests/layer_tests/onnx_tests/test_log.py index db0a329aa09746..53e2c42505bf7b 100644 --- a/tests/layer_tests/onnx_tests/test_log.py +++ b/tests/layer_tests/onnx_tests/test_log.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_logsoftmax.py b/tests/layer_tests/onnx_tests/test_logsoftmax.py index a81b20402d50dd..057376d6ed48b2 100644 --- a/tests/layer_tests/onnx_tests/test_logsoftmax.py +++ b/tests/layer_tests/onnx_tests/test_logsoftmax.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_lrn.py b/tests/layer_tests/onnx_tests/test_lrn.py index 0e8f34129a300f..1c1cf62d5d12b4 100644 --- a/tests/layer_tests/onnx_tests/test_lrn.py +++ b/tests/layer_tests/onnx_tests/test_lrn.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_neg.py b/tests/layer_tests/onnx_tests/test_neg.py index d19991cb8a6b12..98f6acd728f637 100644 --- a/tests/layer_tests/onnx_tests/test_neg.py +++ b/tests/layer_tests/onnx_tests/test_neg.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_non_zero.py b/tests/layer_tests/onnx_tests/test_non_zero.py index 464304651a2a19..a2035b4ab27d63 100644 --- a/tests/layer_tests/onnx_tests/test_non_zero.py +++ b/tests/layer_tests/onnx_tests/test_non_zero.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_not.py b/tests/layer_tests/onnx_tests/test_not.py index 05a6c7ffbb2e2d..1caf8e2e7a770c 100644 --- a/tests/layer_tests/onnx_tests/test_not.py +++ b/tests/layer_tests/onnx_tests/test_not.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_or.py b/tests/layer_tests/onnx_tests/test_or.py index 285c90765d6a7e..6db35aff2f500e 100644 --- a/tests/layer_tests/onnx_tests/test_or.py +++ b/tests/layer_tests/onnx_tests/test_or.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_pad.py b/tests/layer_tests/onnx_tests/test_pad.py index abacc530d93144..161db0685b6fa8 100644 --- a/tests/layer_tests/onnx_tests/test_pad.py +++ b/tests/layer_tests/onnx_tests/test_pad.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_pooling.py b/tests/layer_tests/onnx_tests/test_pooling.py index 85e7fc883fc5d8..2bc2251f8aea49 100644 --- a/tests/layer_tests/onnx_tests/test_pooling.py +++ b/tests/layer_tests/onnx_tests/test_pooling.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_prelu.py b/tests/layer_tests/onnx_tests/test_prelu.py index f20e89b7006a44..59a1e8f4f415e1 100644 --- a/tests/layer_tests/onnx_tests/test_prelu.py +++ b/tests/layer_tests/onnx_tests/test_prelu.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_reduce.py b/tests/layer_tests/onnx_tests/test_reduce.py index 58141e18260016..46b4008c4e653d 100644 --- a/tests/layer_tests/onnx_tests/test_reduce.py +++ b/tests/layer_tests/onnx_tests/test_reduce.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_reduce_lp.py b/tests/layer_tests/onnx_tests/test_reduce_lp.py index 2ff4511ef87443..3cf2f5e133b895 100644 --- a/tests/layer_tests/onnx_tests/test_reduce_lp.py +++ b/tests/layer_tests/onnx_tests/test_reduce_lp.py @@ -5,6 +5,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_relu.py b/tests/layer_tests/onnx_tests/test_relu.py index ce597920923289..520749ed948b25 100644 --- a/tests/layer_tests/onnx_tests/test_relu.py +++ b/tests/layer_tests/onnx_tests/test_relu.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_reshape.py b/tests/layer_tests/onnx_tests/test_reshape.py index 637beeb4388bbb..28eb339af52f9e 100644 --- a/tests/layer_tests/onnx_tests/test_reshape.py +++ b/tests/layer_tests/onnx_tests/test_reshape.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_resize.py b/tests/layer_tests/onnx_tests/test_resize.py index 4d28afdb50fe38..36a808fa859ef1 100644 --- a/tests/layer_tests/onnx_tests/test_resize.py +++ b/tests/layer_tests/onnx_tests/test_resize.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_roi_align.py b/tests/layer_tests/onnx_tests/test_roi_align.py index 4cd49c50c20bf8..d5cedf4e1a0f06 100644 --- a/tests/layer_tests/onnx_tests/test_roi_align.py +++ b/tests/layer_tests/onnx_tests/test_roi_align.py @@ -5,6 +5,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model from unit_tests.utils.graph import build_graph diff --git a/tests/layer_tests/onnx_tests/test_scatter.py b/tests/layer_tests/onnx_tests/test_scatter.py index 578300e144bc3d..baaa0392553fbf 100644 --- a/tests/layer_tests/onnx_tests/test_scatter.py +++ b/tests/layer_tests/onnx_tests/test_scatter.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_sigmoid.py b/tests/layer_tests/onnx_tests/test_sigmoid.py index 5dcb3e8f1b112a..db055a6d9030ac 100644 --- a/tests/layer_tests/onnx_tests/test_sigmoid.py +++ b/tests/layer_tests/onnx_tests/test_sigmoid.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_sign.py b/tests/layer_tests/onnx_tests/test_sign.py index 07f4f169a7bc1b..70c0ffcc0033ec 100644 --- a/tests/layer_tests/onnx_tests/test_sign.py +++ b/tests/layer_tests/onnx_tests/test_sign.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_softmax.py b/tests/layer_tests/onnx_tests/test_softmax.py index c4d9d600276402..390b1a894549c3 100644 --- a/tests/layer_tests/onnx_tests/test_softmax.py +++ b/tests/layer_tests/onnx_tests/test_softmax.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_softplus.py b/tests/layer_tests/onnx_tests/test_softplus.py index cdcbbbf3e8ed13..b0127c0dcf0624 100644 --- a/tests/layer_tests/onnx_tests/test_softplus.py +++ b/tests/layer_tests/onnx_tests/test_softplus.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_softsign.py b/tests/layer_tests/onnx_tests/test_softsign.py index 30ca27402c7878..75043b57b80dc7 100644 --- a/tests/layer_tests/onnx_tests/test_softsign.py +++ b/tests/layer_tests/onnx_tests/test_softsign.py @@ -2,6 +2,8 @@ # SPDX-License-Identifier: Apache-2.0 import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_sqrt.py b/tests/layer_tests/onnx_tests/test_sqrt.py index 9c4733a68cd9fa..24dbbcac659df4 100644 --- a/tests/layer_tests/onnx_tests/test_sqrt.py +++ b/tests/layer_tests/onnx_tests/test_sqrt.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_trigonometry.py b/tests/layer_tests/onnx_tests/test_trigonometry.py index 563b63b1e5632d..99651091ea2e96 100644 --- a/tests/layer_tests/onnx_tests/test_trigonometry.py +++ b/tests/layer_tests/onnx_tests/test_trigonometry.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_where.py b/tests/layer_tests/onnx_tests/test_where.py index fb358a2ced8415..1bf845340b3922 100644 --- a/tests/layer_tests/onnx_tests/test_where.py +++ b/tests/layer_tests/onnx_tests/test_where.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/onnx_tests/test_xor.py b/tests/layer_tests/onnx_tests/test_xor.py index 2790a31784ff59..e7f0c11f8362a2 100644 --- a/tests/layer_tests/onnx_tests/test_xor.py +++ b/tests/layer_tests/onnx_tests/test_xor.py @@ -3,6 +3,8 @@ import numpy as np import pytest +pytest.importorskip("openvino.tools.mo", reason="Ticket - 157136") + from common.layer_test_class import check_ir_version from common.onnx_layer_test_class import OnnxRuntimeLayerTest, onnx_make_model diff --git a/tests/layer_tests/requirements.txt b/tests/layer_tests/requirements.txt index 04889ebce10a39..2ba12cc5e2bece 100644 --- a/tests/layer_tests/requirements.txt +++ b/tests/layer_tests/requirements.txt @@ -16,5 +16,3 @@ pytest defusedxml tensorflow tensorflow-addons; python_version <= '3.10' -jax; sys_platform == "linux" and platform_machine == "x86_64" # https://jax.readthedocs.io/en/latest/installation.html#pip-installation-cpu - wheels are for "x86_64" only -jaxlib; sys_platform == "linux" and platform_machine == "x86_64" # https://jax.readthedocs.io/en/latest/installation.html#pip-installation-cpu - wheels are for "x86_64" only diff --git a/tests/requirements_pytorch b/tests/requirements_pytorch index f42deb81839883..33907145f7de4b 100644 --- a/tests/requirements_pytorch +++ b/tests/requirements_pytorch @@ -14,7 +14,8 @@ torchaudio==2.2.2; platform_system == "Darwin" and platform_machine == "x86_64" # transformers 4.45.1 is available # but optimum still requires <4.45.0 transformers==4.44.2 -pytest==7.0.1 +pytest==7.0.1; python_version < '3.10' +pytest==7.2.0; python_version >= '3.10' pytest-html==4.1.1 pytest-xdist[psutil]==3.6.1 defusedxml==0.7.1 diff --git a/tests/requirements_tensorflow b/tests/requirements_tensorflow index 8e0d1141695ef9..5d699facad1c91 100644 --- a/tests/requirements_tensorflow +++ b/tests/requirements_tensorflow @@ -4,7 +4,8 @@ # tensorflow 2.16.2 depends on numpy<2.0.0 and >=1.26.0; python_version >= "3.12" numpy==1.26.4; python_version < "3.12" or platform_system == "Darwin" and platform_machine == "x86_64" numpy==2.0.2; python_version >= "3.12" and (platform_system != "Darwin" or platform_machine != "x86_64") -pytest==7.0.1 +pytest==7.0.1; python_version < '3.10' +pytest==7.2.0; python_version >= '3.10' pytest-xdist[psutil]==3.6.1 pytest-html==4.1.1 transformers==4.45.1 diff --git a/tools/benchmark_tool/openvino/__init__.py b/tools/benchmark_tool/openvino/__init__.py index 7643f742e0067d..69c678909b1c9e 100644 --- a/tools/benchmark_tool/openvino/__init__.py +++ b/tools/benchmark_tool/openvino/__init__.py @@ -7,7 +7,7 @@ # Required for Windows OS platforms # Note: always top-level try: - from openvino.package_utils import _add_openvino_libs_to_search_path + from openvino.utils import _add_openvino_libs_to_search_path _add_openvino_libs_to_search_path() except ImportError: pass @@ -17,47 +17,6 @@ # # This __init__.py forces checking of runtime modules to propagate errors. # # It is not compared with init files from openvino-dev package. # # - -# Openvino pybind bindings -from openvino._pyopenvino import AxisSet -from openvino._pyopenvino import AxisVector -from openvino._pyopenvino import ConstOutput -from openvino._pyopenvino import Coordinate -from openvino._pyopenvino import CoordinateDiff -from openvino._pyopenvino import DiscreteTypeInfo -from openvino._pyopenvino import Extension -from openvino._pyopenvino import ProfilingInfo -from openvino._pyopenvino import RTMap -from openvino._pyopenvino import Version -from openvino._pyopenvino import Symbol -from openvino._pyopenvino import Dimension -from openvino._pyopenvino import Input -from openvino._pyopenvino import Output -from openvino._pyopenvino import Node -from openvino._pyopenvino import Strides -from openvino._pyopenvino import PartialShape -from openvino._pyopenvino import Shape -from openvino._pyopenvino import Layout -from openvino._pyopenvino import Type -from openvino._pyopenvino import Tensor -from openvino._pyopenvino import OVAny -from openvino._pyopenvino import get_batch -from openvino._pyopenvino import set_batch -from openvino._pyopenvino import serialize -from openvino._pyopenvino import shutdown -from openvino._pyopenvino import save_model -from openvino._pyopenvino import layout_helpers -from openvino._pyopenvino import RemoteContext -from openvino._pyopenvino import RemoteTensor -from openvino._pyopenvino import Op - -# Import public classes from _ov_api -from openvino._ov_api import Model -from openvino._ov_api import Core -from openvino._ov_api import CompiledModel -from openvino._ov_api import InferRequest -from openvino._ov_api import AsyncInferQueue - # Import all public modules from openvino import runtime as runtime from openvino import frontend as frontend @@ -67,10 +26,36 @@ from openvino import utils as utils from openvino import properties as properties +# Import most important classes and functions from openvino.runtime +from openvino._ov_api import Model +from openvino._ov_api import Core +from openvino._ov_api import CompiledModel +from openvino._ov_api import InferRequest +from openvino._ov_api import AsyncInferQueue + +from openvino.runtime import Symbol +from openvino.runtime import Dimension +from openvino.runtime import Strides +from openvino.runtime import PartialShape +from openvino.runtime import Shape +from openvino.runtime import Layout +from openvino.runtime import Type +from openvino.runtime import Tensor +from openvino.runtime import OVAny + # Helper functions for openvino module -from openvino.utils.data_helpers import tensor_from_file +from openvino.runtime.utils.data_helpers import tensor_from_file from openvino._ov_api import compile_model +from openvino.runtime import get_batch +from openvino.runtime import set_batch +from openvino.runtime import serialize +from openvino.runtime import shutdown +from openvino.runtime import save_model +from openvino.runtime import layout_helpers +from openvino._pyopenvino import RemoteContext +from openvino._pyopenvino import RemoteTensor +from openvino._pyopenvino import Op # Import opsets from openvino import opset1 @@ -95,7 +80,7 @@ from openvino._pyopenvino import VASurfaceTensor # Set version for openvino package -from openvino._pyopenvino import get_version +from openvino.runtime import get_version __version__ = get_version() # Tools diff --git a/tools/mo/openvino/__init__.py b/tools/mo/openvino/__init__.py index 7643f742e0067d..b015570964c520 100644 --- a/tools/mo/openvino/__init__.py +++ b/tools/mo/openvino/__init__.py @@ -7,96 +7,61 @@ # Required for Windows OS platforms # Note: always top-level try: - from openvino.package_utils import _add_openvino_libs_to_search_path + from openvino.utils import _add_openvino_libs_to_search_path _add_openvino_libs_to_search_path() except ImportError: pass -# # -# # OpenVINO API -# # This __init__.py forces checking of runtime modules to propagate errors. -# # It is not compared with init files from openvino-dev package. -# # - -# Openvino pybind bindings -from openvino._pyopenvino import AxisSet -from openvino._pyopenvino import AxisVector -from openvino._pyopenvino import ConstOutput -from openvino._pyopenvino import Coordinate -from openvino._pyopenvino import CoordinateDiff -from openvino._pyopenvino import DiscreteTypeInfo -from openvino._pyopenvino import Extension -from openvino._pyopenvino import ProfilingInfo -from openvino._pyopenvino import RTMap -from openvino._pyopenvino import Version -from openvino._pyopenvino import Symbol -from openvino._pyopenvino import Dimension -from openvino._pyopenvino import Input -from openvino._pyopenvino import Output -from openvino._pyopenvino import Node -from openvino._pyopenvino import Strides -from openvino._pyopenvino import PartialShape -from openvino._pyopenvino import Shape -from openvino._pyopenvino import Layout -from openvino._pyopenvino import Type -from openvino._pyopenvino import Tensor -from openvino._pyopenvino import OVAny -from openvino._pyopenvino import get_batch -from openvino._pyopenvino import set_batch -from openvino._pyopenvino import serialize -from openvino._pyopenvino import shutdown -from openvino._pyopenvino import save_model -from openvino._pyopenvino import layout_helpers -from openvino._pyopenvino import RemoteContext -from openvino._pyopenvino import RemoteTensor -from openvino._pyopenvino import Op - -# Import public classes from _ov_api -from openvino._ov_api import Model -from openvino._ov_api import Core -from openvino._ov_api import CompiledModel -from openvino._ov_api import InferRequest -from openvino._ov_api import AsyncInferQueue +# OpenVINO API +try: + # Import all public modules + from openvino import runtime as runtime + from openvino import frontend as frontend + from openvino import helpers as helpers + from openvino import preprocess as preprocess + from openvino import utils as utils + from openvino import properties as properties -# Import all public modules -from openvino import runtime as runtime -from openvino import frontend as frontend -from openvino import helpers as helpers -from openvino import experimental as experimental -from openvino import preprocess as preprocess -from openvino import utils as utils -from openvino import properties as properties + # Import most important classes and functions from openvino.runtime + from openvino.runtime import Model + from openvino.runtime import Core + from openvino.runtime import CompiledModel + from openvino.runtime import InferRequest + from openvino.runtime import AsyncInferQueue -# Helper functions for openvino module -from openvino.utils.data_helpers import tensor_from_file -from openvino._ov_api import compile_model + from openvino.runtime import Symbol + from openvino.runtime import Dimension + from openvino.runtime import Strides + from openvino.runtime import PartialShape + from openvino.runtime import Shape + from openvino.runtime import Layout + from openvino.runtime import Type + from openvino.runtime import Tensor + from openvino.runtime import OVAny + from openvino.runtime import compile_model + from openvino.runtime import get_batch + from openvino.runtime import set_batch + from openvino.runtime import serialize + from openvino.runtime import shutdown + from openvino.runtime import tensor_from_file + from openvino.runtime import save_model + from openvino.runtime import layout_helpers -# Import opsets -from openvino import opset1 -from openvino import opset2 -from openvino import opset3 -from openvino import opset4 -from openvino import opset5 -from openvino import opset6 -from openvino import opset7 -from openvino import opset8 -from openvino import opset9 -from openvino import opset10 -from openvino import opset11 -from openvino import opset12 -from openvino import opset13 -from openvino import opset14 -from openvino import opset15 -from openvino import opset16 + from openvino._pyopenvino import RemoteContext + from openvino._pyopenvino import RemoteTensor + from openvino._pyopenvino import Op -# libva related: -from openvino._pyopenvino import VAContext -from openvino._pyopenvino import VASurfaceTensor + # libva related: + from openvino._pyopenvino import VAContext + from openvino._pyopenvino import VASurfaceTensor -# Set version for openvino package -from openvino._pyopenvino import get_version -__version__ = get_version() + # Set version for openvino package + from openvino.runtime import get_version + __version__ = get_version() +except ImportError: + import warnings + warnings.warn("openvino package has problems with imports!", ImportWarning, stacklevel=2) # Tools try: diff --git a/tools/openvino_dev/src/openvino/__init__.py b/tools/openvino_dev/src/openvino/__init__.py index 7643f742e0067d..b015570964c520 100644 --- a/tools/openvino_dev/src/openvino/__init__.py +++ b/tools/openvino_dev/src/openvino/__init__.py @@ -7,96 +7,61 @@ # Required for Windows OS platforms # Note: always top-level try: - from openvino.package_utils import _add_openvino_libs_to_search_path + from openvino.utils import _add_openvino_libs_to_search_path _add_openvino_libs_to_search_path() except ImportError: pass -# # -# # OpenVINO API -# # This __init__.py forces checking of runtime modules to propagate errors. -# # It is not compared with init files from openvino-dev package. -# # - -# Openvino pybind bindings -from openvino._pyopenvino import AxisSet -from openvino._pyopenvino import AxisVector -from openvino._pyopenvino import ConstOutput -from openvino._pyopenvino import Coordinate -from openvino._pyopenvino import CoordinateDiff -from openvino._pyopenvino import DiscreteTypeInfo -from openvino._pyopenvino import Extension -from openvino._pyopenvino import ProfilingInfo -from openvino._pyopenvino import RTMap -from openvino._pyopenvino import Version -from openvino._pyopenvino import Symbol -from openvino._pyopenvino import Dimension -from openvino._pyopenvino import Input -from openvino._pyopenvino import Output -from openvino._pyopenvino import Node -from openvino._pyopenvino import Strides -from openvino._pyopenvino import PartialShape -from openvino._pyopenvino import Shape -from openvino._pyopenvino import Layout -from openvino._pyopenvino import Type -from openvino._pyopenvino import Tensor -from openvino._pyopenvino import OVAny -from openvino._pyopenvino import get_batch -from openvino._pyopenvino import set_batch -from openvino._pyopenvino import serialize -from openvino._pyopenvino import shutdown -from openvino._pyopenvino import save_model -from openvino._pyopenvino import layout_helpers -from openvino._pyopenvino import RemoteContext -from openvino._pyopenvino import RemoteTensor -from openvino._pyopenvino import Op - -# Import public classes from _ov_api -from openvino._ov_api import Model -from openvino._ov_api import Core -from openvino._ov_api import CompiledModel -from openvino._ov_api import InferRequest -from openvino._ov_api import AsyncInferQueue +# OpenVINO API +try: + # Import all public modules + from openvino import runtime as runtime + from openvino import frontend as frontend + from openvino import helpers as helpers + from openvino import preprocess as preprocess + from openvino import utils as utils + from openvino import properties as properties -# Import all public modules -from openvino import runtime as runtime -from openvino import frontend as frontend -from openvino import helpers as helpers -from openvino import experimental as experimental -from openvino import preprocess as preprocess -from openvino import utils as utils -from openvino import properties as properties + # Import most important classes and functions from openvino.runtime + from openvino.runtime import Model + from openvino.runtime import Core + from openvino.runtime import CompiledModel + from openvino.runtime import InferRequest + from openvino.runtime import AsyncInferQueue -# Helper functions for openvino module -from openvino.utils.data_helpers import tensor_from_file -from openvino._ov_api import compile_model + from openvino.runtime import Symbol + from openvino.runtime import Dimension + from openvino.runtime import Strides + from openvino.runtime import PartialShape + from openvino.runtime import Shape + from openvino.runtime import Layout + from openvino.runtime import Type + from openvino.runtime import Tensor + from openvino.runtime import OVAny + from openvino.runtime import compile_model + from openvino.runtime import get_batch + from openvino.runtime import set_batch + from openvino.runtime import serialize + from openvino.runtime import shutdown + from openvino.runtime import tensor_from_file + from openvino.runtime import save_model + from openvino.runtime import layout_helpers -# Import opsets -from openvino import opset1 -from openvino import opset2 -from openvino import opset3 -from openvino import opset4 -from openvino import opset5 -from openvino import opset6 -from openvino import opset7 -from openvino import opset8 -from openvino import opset9 -from openvino import opset10 -from openvino import opset11 -from openvino import opset12 -from openvino import opset13 -from openvino import opset14 -from openvino import opset15 -from openvino import opset16 + from openvino._pyopenvino import RemoteContext + from openvino._pyopenvino import RemoteTensor + from openvino._pyopenvino import Op -# libva related: -from openvino._pyopenvino import VAContext -from openvino._pyopenvino import VASurfaceTensor + # libva related: + from openvino._pyopenvino import VAContext + from openvino._pyopenvino import VASurfaceTensor -# Set version for openvino package -from openvino._pyopenvino import get_version -__version__ = get_version() + # Set version for openvino package + from openvino.runtime import get_version + __version__ = get_version() +except ImportError: + import warnings + warnings.warn("openvino package has problems with imports!", ImportWarning, stacklevel=2) # Tools try: diff --git a/tools/ovc/openvino/__init__.py b/tools/ovc/openvino/__init__.py index 7643f742e0067d..69c678909b1c9e 100644 --- a/tools/ovc/openvino/__init__.py +++ b/tools/ovc/openvino/__init__.py @@ -7,7 +7,7 @@ # Required for Windows OS platforms # Note: always top-level try: - from openvino.package_utils import _add_openvino_libs_to_search_path + from openvino.utils import _add_openvino_libs_to_search_path _add_openvino_libs_to_search_path() except ImportError: pass @@ -17,47 +17,6 @@ # # This __init__.py forces checking of runtime modules to propagate errors. # # It is not compared with init files from openvino-dev package. # # - -# Openvino pybind bindings -from openvino._pyopenvino import AxisSet -from openvino._pyopenvino import AxisVector -from openvino._pyopenvino import ConstOutput -from openvino._pyopenvino import Coordinate -from openvino._pyopenvino import CoordinateDiff -from openvino._pyopenvino import DiscreteTypeInfo -from openvino._pyopenvino import Extension -from openvino._pyopenvino import ProfilingInfo -from openvino._pyopenvino import RTMap -from openvino._pyopenvino import Version -from openvino._pyopenvino import Symbol -from openvino._pyopenvino import Dimension -from openvino._pyopenvino import Input -from openvino._pyopenvino import Output -from openvino._pyopenvino import Node -from openvino._pyopenvino import Strides -from openvino._pyopenvino import PartialShape -from openvino._pyopenvino import Shape -from openvino._pyopenvino import Layout -from openvino._pyopenvino import Type -from openvino._pyopenvino import Tensor -from openvino._pyopenvino import OVAny -from openvino._pyopenvino import get_batch -from openvino._pyopenvino import set_batch -from openvino._pyopenvino import serialize -from openvino._pyopenvino import shutdown -from openvino._pyopenvino import save_model -from openvino._pyopenvino import layout_helpers -from openvino._pyopenvino import RemoteContext -from openvino._pyopenvino import RemoteTensor -from openvino._pyopenvino import Op - -# Import public classes from _ov_api -from openvino._ov_api import Model -from openvino._ov_api import Core -from openvino._ov_api import CompiledModel -from openvino._ov_api import InferRequest -from openvino._ov_api import AsyncInferQueue - # Import all public modules from openvino import runtime as runtime from openvino import frontend as frontend @@ -67,10 +26,36 @@ from openvino import utils as utils from openvino import properties as properties +# Import most important classes and functions from openvino.runtime +from openvino._ov_api import Model +from openvino._ov_api import Core +from openvino._ov_api import CompiledModel +from openvino._ov_api import InferRequest +from openvino._ov_api import AsyncInferQueue + +from openvino.runtime import Symbol +from openvino.runtime import Dimension +from openvino.runtime import Strides +from openvino.runtime import PartialShape +from openvino.runtime import Shape +from openvino.runtime import Layout +from openvino.runtime import Type +from openvino.runtime import Tensor +from openvino.runtime import OVAny + # Helper functions for openvino module -from openvino.utils.data_helpers import tensor_from_file +from openvino.runtime.utils.data_helpers import tensor_from_file from openvino._ov_api import compile_model +from openvino.runtime import get_batch +from openvino.runtime import set_batch +from openvino.runtime import serialize +from openvino.runtime import shutdown +from openvino.runtime import save_model +from openvino.runtime import layout_helpers +from openvino._pyopenvino import RemoteContext +from openvino._pyopenvino import RemoteTensor +from openvino._pyopenvino import Op # Import opsets from openvino import opset1 @@ -95,7 +80,7 @@ from openvino._pyopenvino import VASurfaceTensor # Set version for openvino package -from openvino._pyopenvino import get_version +from openvino.runtime import get_version __version__ = get_version() # Tools