diff --git a/README.md b/README.md index 695f84f1628118..7e9b173530de61 100644 --- a/README.md +++ b/README.md @@ -1,6 +1,14 @@
+

+Open-source software toolkit for optimizing and deploying deep learning models. +

+ +

+ Documentation • Blog • Key Features • Tutorials • Integrations • Benchmarks • Generative AI +

+ [![PyPI Status](https://badge.fury.io/py/openvino.svg)](https://badge.fury.io/py/openvino) [![Anaconda Status](https://anaconda.org/conda-forge/openvino/badges/version.svg)](https://anaconda.org/conda-forge/openvino) [![brew Status](https://img.shields.io/homebrew/v/openvino)](https://formulae.brew.sh/formula/openvino) @@ -10,14 +18,14 @@ [![brew Downloads](https://img.shields.io/homebrew/installs/dy/openvino)](https://formulae.brew.sh/formula/openvino)
-Welcome to OpenVINO™, an open-source software toolkit for optimizing and deploying deep learning models. - **Inference Optimization**: Boost deep learning performance in computer vision, automatic speech recognition, generative AI, natural language processing with large and small language models, and many other common tasks. -- **Flexible Model Support**: Use models trained with popular frameworks such as TensorFlow, PyTorch, ONNX, Keras, and PaddlePaddle. Convert and deploy models without original frameworks. +- **Flexible Model Support**: Use models trained with popular frameworks such as PyTorch, TensorFlow, ONNX, Keras, PaddlePaddle, and JAX/Flax. Directly integrate models built with transformers and diffusers from the Hugging Face Hub using Optimum Intel. Convert and deploy models without original frameworks. - **Broad Platform Compatibility**: Reduce resource demands and efficiently deploy on a range of platforms from edge to cloud. OpenVINO™ supports inference on CPU (x86, ARM), GPU (OpenCL capable, integrated and discrete) and AI accelerators (Intel NPU). - **Community and Ecosystem**: Join an active community contributing to the enhancement of deep learning performance across various domains. -Check out the [OpenVINO Cheat Sheet](https://docs.openvino.ai/2024/_static/download/OpenVINO_Quick_Start_Guide.pdf) for a quick reference. +Check out the [OpenVINO Cheat Sheet](https://docs.openvino.ai/2024/_static/download/OpenVINO_Quick_Start_Guide.pdf) and [Key Features](https://docs.openvino.ai/2024/about-openvino/key-features.html) for a quick reference. + ## Installation @@ -40,6 +48,8 @@ Learn how to optimize and deploy popular models with the [OpenVINO Notebooks](ht - [Multimodal assistant with LLaVa and OpenVINO](https://github.com/openvinotoolkit/openvino_notebooks/blob/latest/notebooks/llava-multimodal-chatbot/llava-multimodal-chatbot-genai.ipynb) - [Automatic speech recognition using Whisper and OpenVINO](https://github.com/openvinotoolkit/openvino_notebooks/blob/latest/notebooks/whisper-asr-genai/whisper-asr-genai.ipynb) +Discover more examples in the [OpenVINO Samples (Python & C++)](https://docs.openvino.ai/2024/learn-openvino/openvino-samples.html) and [Notebooks (Python)](https://docs.openvino.ai/2024/learn-openvino/interactive-tutorials-python.html). + Here are easy-to-follow code examples demonstrating how to run PyTorch and TensorFlow model inference using OpenVINO: **PyTorch Model** @@ -86,25 +96,43 @@ data = np.random.rand(1, 224, 224, 3) output = compiled_model({0: data}) ``` -OpenVINO also supports CPU, GPU, and NPU devices and works with models in TensorFlow, PyTorch, ONNX, TensorFlow Lite, PaddlePaddle model formats. -With OpenVINO you can do automatic performance enhancements at runtime customized to your hardware (preserving model accuracy), including: -asynchronous execution, batch processing, tensor fusion, load balancing, dynamic inference parallelism, automatic BF16 conversion, and more. +OpenVINO supports the CPU, GPU, and NPU [devices](https://docs.openvino.ai/2024/openvino-workflow/running-inference/inference-devices-and-modes.html) and works with models from PyTorch, TensorFlow, ONNX, TensorFlow Lite, PaddlePaddle, and JAX/Flax [frameworks](https://docs.openvino.ai/2024/openvino-workflow/model-preparation.html). It includes [APIs](https://docs.openvino.ai/2024/api/api_reference.html) in C++, Python, C, NodeJS, and offers the GenAI API for optimized model pipelines and performance. + +## Generative AI with OpenVINO + +Get started with the OpenVINO GenAI [installation](https://docs.openvino.ai/2024/get-started/install-openvino/install-openvino-genai.html) and refer to the [detailed guide](https://docs.openvino.ai/2024/learn-openvino/llm_inference_guide/genai-guide.html) to explore the capabilities of Generative AI using OpenVINO. + +Learn how to run LLMs and GenAI with [Samples](https://github.com/openvinotoolkit/openvino.genai/tree/master/samples) in the [OpenVINO™ GenAI repo](https://github.com/openvinotoolkit/openvino.genai). See GenAI in action with Jupyter notebooks: [LLM-powered Chatbot](https://github.com/openvinotoolkit/openvino_notebooks/blob/latest/notebooks/llm-chatbot/README.md) and [LLM Instruction-following pipeline](https://github.com/openvinotoolkit/openvino_notebooks/blob/latest/notebooks/llm-question-answering/README.md). + +## Documentation + +[User documentation](https://docs.openvino.ai/) contains detailed information about OpenVINO and guides you from installation through optimizing and deploying models for your AI applications. + +[Developer documentation](./docs/dev/index.md) focuses on the OpenVINO architecture and describes [building](./docs/dev/build.md) and [contributing](./CONTRIBUTING.md) processes. ## OpenVINO Ecosystem -- [🤗Optimum Intel](https://github.com/huggingface/optimum-intel) - a simple interface to optimize Transformers and Diffusers models. +### OpenVINO Tools + - [Neural Network Compression Framework (NNCF)](https://github.com/openvinotoolkit/nncf) - advanced model optimization techniques including quantization, filter pruning, binarization, and sparsity. - [GenAI Repository](https://github.com/openvinotoolkit/openvino.genai) and [OpenVINO Tokenizers](https://github.com/openvinotoolkit/openvino_tokenizers) - resources and tools for developing and optimizing Generative AI applications. - [OpenVINO™ Model Server (OVMS)](https://github.com/openvinotoolkit/model_server) - a scalable, high-performance solution for serving models optimized for Intel architectures. - [Intel® Geti™](https://geti.intel.com/) - an interactive video and image annotation tool for computer vision use cases. -Check out the [Awesome OpenVINO](https://github.com/openvinotoolkit/awesome-openvino) repository to discover a collection of community-made AI projects based on OpenVINO! +### Integrations -## Documentation +- [🤗Optimum Intel](https://github.com/huggingface/optimum-intel) - grab and use models leveraging OpenVINO within the Hugging Face API. +- [Torch.compile](https://docs.openvino.ai/2024/openvino-workflow/torch-compile.html) - use OpenVINO for Python-native applications by JIT-compiling code into optimized kernels. +- [OpenVINO LLMs inference and serving with vLLM​](https://docs.vllm.ai/en/stable/getting_started/openvino-installation.html) - enhance vLLM's fast and easy model serving with the OpenVINO backend. +- [OpenVINO Execution Provider for ONNX Runtime](https://onnxruntime.ai/docs/execution-providers/OpenVINO-ExecutionProvider.html) - use OpenVINO as a backend with your existing ONNX Runtime code. +- [LlamaIndex](https://docs.llamaindex.ai/en/stable/examples/llm/openvino/) - build context-augmented GenAI applications with the LlamaIndex framework and enhance runtime performance with OpenVINO. +- [LangChain](https://python.langchain.com/docs/integrations/llms/openvino/) - integrate OpenVINO with the LangChain framework to enhance runtime performance for GenAI applications. -[User documentation](https://docs.openvino.ai/) contains detailed information about OpenVINO and guides you from installation through optimizing and deploying models for your AI applications. +Check out the [Awesome OpenVINO](https://github.com/openvinotoolkit/awesome-openvino) repository to discover a collection of community-made AI projects based on OpenVINO! -[Developer documentation](./docs/dev/index.md) focuses on how OpenVINO [components](./docs/dev/index.md#openvino-components) work and describes [building](./docs/dev/build.md) and [contributing](./CONTRIBUTING.md) processes. +## Performance + +Explore [OpenVINO Performance Benchmarks](https://docs.openvino.ai/2024/about-openvino/performance-benchmarks.html) to discover the optimal hardware configurations and plan your AI deployment based on verified data. ## Contribution and Support @@ -118,9 +146,8 @@ You can ask questions and get support on: * The [`openvino`](https://stackoverflow.com/questions/tagged/openvino) tag on Stack Overflow\*. -## Additional Resources +## Resources -* [Product Page](https://software.intel.com/content/www/us/en/develop/tools/openvino-toolkit.html) * [Release Notes](https://docs.openvino.ai/2024/about-openvino/release-notes-openvino.html) * [OpenVINO Blog](https://blog.openvino.ai/) * [OpenVINO™ toolkit on Medium](https://medium.com/@openvino) @@ -145,4 +172,3 @@ By contributing to the project, you agree to the license and copyright terms the --- \* Other names and brands may be claimed as the property of others. - diff --git a/cmake/developer_package/cross_compile/cross_compiled_func.cmake b/cmake/developer_package/cross_compile/cross_compiled_func.cmake index 962aa5d373a4db..d83450e6d238bd 100644 --- a/cmake/developer_package/cross_compile/cross_compiled_func.cmake +++ b/cmake/developer_package/cross_compile/cross_compiled_func.cmake @@ -11,7 +11,7 @@ set(_ACCEPTED_ARCHS_AVX "^(ANY|SSE42|AVX)$") set(_ACCEPTED_ARCHS_AVX2 "^(ANY|SSE42|AVX|AVX2)$") set(_ACCEPTED_ARCHS_AVX512F "^(ANY|SSE42|AVX|AVX2|AVX512F)$") set(_ACCEPTED_ARCHS_NEON_FP16 "^(ANY|NEON_FP16)$") -set(_ACCEPTED_ARCHS_SVE "^(ANY|SVE)$") +set(_ACCEPTED_ARCHS_SVE "^(ANY|NEON_FP16|SVE)$") ## Arch specific definitions set(_DEFINE_ANY "") @@ -186,10 +186,10 @@ endfunction() # Return currently requested ARCH id # function(_currently_requested_top_arch VAR) - if(ENABLE_NEON_FP16) - set(RES NEON_FP16) - elseif(ENABLE_SVE) + if(ENABLE_SVE) set(RES SVE) + elseif(ENABLE_NEON_FP16) + set(RES NEON_FP16) elseif(ENABLE_AVX512F) set(RES AVX512F) elseif(ENABLE_AVX2) diff --git a/cmake/developer_package/features.cmake b/cmake/developer_package/features.cmake index ae5313cea8a8b4..538ec2e4da1c79 100644 --- a/cmake/developer_package/features.cmake +++ b/cmake/developer_package/features.cmake @@ -4,6 +4,7 @@ include(options) include(target_flags) +include(compile_flags/os_flags) if(WIN32) set (CPACK_GENERATOR "ZIP" CACHE STRING "Cpack generator for OpenVINO") @@ -49,9 +50,9 @@ ov_dependent_option (ENABLE_AVX2 "Enable AVX2 optimizations" ON "X86_64 OR (X86 ov_dependent_option (ENABLE_AVX512F "Enable AVX512 optimizations" ON "X86_64 OR (X86 AND NOT EMSCRIPTEN)" OFF) -ov_dependent_option(ENABLE_NEON_FP16 "Enable ARM FP16 optimizations" ON "AARCH64" OFF) +ov_dependent_option (ENABLE_NEON_FP16 "Enable ARM FP16 optimizations" ON "AARCH64" OFF) -ov_dependent_option(ENABLE_SVE "Enable SVE optimizations" ON "AARCH64" OFF) +ov_dependent_option (ENABLE_SVE "Enable SVE optimizations" ON "AARCH64" OFF) # Type of build, we add this as an explicit option to default it to ON get_property(BUILD_SHARED_LIBS_DEFAULT GLOBAL PROPERTY TARGET_SUPPORTS_SHARED_LIBS) @@ -106,3 +107,11 @@ if(ENABLE_AVX512F) set(ENABLE_AVX512F OFF CACHE BOOL "" FORCE) endif() endif() + +if(ENABLE_SVE) + ov_check_compiler_supports_sve("-march=armv8-a+sve") + + if(NOT CXX_HAS_SVE) + set(ENABLE_SVE OFF CACHE BOOL "" FORCE) + endif() +endif() \ No newline at end of file diff --git a/docs/articles_en/about-openvino/performance-benchmarks.rst b/docs/articles_en/about-openvino/performance-benchmarks.rst index a408253038e75c..4262ec6b2b3732 100644 --- a/docs/articles_en/about-openvino/performance-benchmarks.rst +++ b/docs/articles_en/about-openvino/performance-benchmarks.rst @@ -56,8 +56,7 @@ implemented in your solutions. Click the buttons below to see the chosen benchma :material-regular:`table_view;1.4em` LLM performance for AI PC -.. uncomment under - .. .. grid-item:: + .. grid-item:: .. button-link:: # :class: ovms-toolkit-benchmark-llm-result diff --git a/docs/articles_en/about-openvino/release-notes-openvino.rst b/docs/articles_en/about-openvino/release-notes-openvino.rst index cdce4fe86f2aaa..bf475159380dff 100644 --- a/docs/articles_en/about-openvino/release-notes-openvino.rst +++ b/docs/articles_en/about-openvino/release-notes-openvino.rst @@ -641,7 +641,7 @@ Previous 2024 releases * New samples and pipelines are now available: * An example IterableStreamer implementation in - `multinomial_causal_lm/python sample `__ + `multinomial_causal_lm/python sample `__ * GenAI compilation is now available as part of OpenVINO via the –DOPENVINO_EXTRA_MODULES CMake option. diff --git a/docs/articles_en/learn-openvino/llm_inference_guide/genai-guide.rst b/docs/articles_en/learn-openvino/llm_inference_guide/genai-guide.rst index dbc5d3c4416cd4..43f9435bf79b1b 100644 --- a/docs/articles_en/learn-openvino/llm_inference_guide/genai-guide.rst +++ b/docs/articles_en/learn-openvino/llm_inference_guide/genai-guide.rst @@ -367,7 +367,7 @@ make sure to :doc:`install OpenVINO with GenAI <../../get-started/install-openvi For more information, refer to the - `Python sample `__. + `Python sample `__. .. tab-item:: C++ :sync: cpp @@ -415,7 +415,7 @@ make sure to :doc:`install OpenVINO with GenAI <../../get-started/install-openvi For more information, refer to the - `C++ sample `__ + `C++ sample `__ .. dropdown:: Using GenAI with Vision Language Models @@ -803,7 +803,7 @@ runs prediction of the next K tokens, thus repeating the cycle. For more information, refer to the - `Python sample `__. + `Python sample `__. .. tab-item:: C++ @@ -859,7 +859,7 @@ runs prediction of the next K tokens, thus repeating the cycle. For more information, refer to the - `C++ sample `__ + `C++ sample `__ diff --git a/docs/articles_en/openvino-workflow/model-preparation/convert-model-pytorch.rst b/docs/articles_en/openvino-workflow/model-preparation/convert-model-pytorch.rst index 62cfdf05f2b11f..fc2637aba9139e 100644 --- a/docs/articles_en/openvino-workflow/model-preparation/convert-model-pytorch.rst +++ b/docs/articles_en/openvino-workflow/model-preparation/convert-model-pytorch.rst @@ -206,14 +206,16 @@ Here is an example of how to convert a model obtained with ``torch.export``: Converting a PyTorch Model from Disk #################################### -PyTorch provides the capability to save models in two distinct formats: ``torch.jit.ScriptModule`` and ``torch.export.ExportedProgram``. -Both formats can be saved to disk as standalone files, enabling them to be reloaded independently of the original Python code. +PyTorch can save models in two formats: ``torch.jit.ScriptModule`` and ``torch.export.ExportedProgram``. +Both formats may be saved to drive as standalone files and reloaded later, independently of the +original Python code. ExportedProgram Format ++++++++++++++++++++++ -The ``ExportedProgram`` format is saved on disk using `torch.export.save() `__. -Below is an example of how to convert an ``ExportedProgram`` from disk: +You can save the ``ExportedProgram`` format using +`torch.export.save() `__. +Here is an example of how to convert it: .. tab-set:: @@ -236,8 +238,9 @@ Below is an example of how to convert an ``ExportedProgram`` from disk: ScriptModule Format +++++++++++++++++++ -`torch.jit.save() `__ serializes ``ScriptModule`` object on disk. -To convert the serialized ``ScriptModule`` format, run ``convert_model`` function with ``example_input`` parameter as follows: +`torch.jit.save() `__ serializes +the ``ScriptModule`` object on a drive. To convert the serialized ``ScriptModule`` format, run +the ``convert_model`` function with ``example_input`` parameter as follows: .. code-block:: py :force: @@ -252,15 +255,15 @@ To convert the serialized ``ScriptModule`` format, run ``convert_model`` functio Exporting a PyTorch Model to ONNX Format ######################################## -An alternative method of converting PyTorch models is exporting a PyTorch model to ONNX with -``torch.onnx.export`` first and then converting the resulting ``.onnx`` file to OpenVINO Model -with ``openvino.convert_model``. It can be considered as a backup solution if a model cannot be -converted directly from PyTorch to OpenVINO as described in the above chapters. Converting through -ONNX can be more expensive in terms of code, conversion time, and allocated memory. +An alternative method of converting a PyTorch models is to export it to ONNX first +(with ``torch.onnx.export``) and then convert the resulting ``.onnx`` file to the OpenVINO IR +model (with ``openvino.convert_model``). It should be considered a backup solution, if a model +cannot be converted directly, as described previously. Converting through ONNX can be more +expensive in terms of code overhead, conversion time, and allocated memory. 1. Refer to the `Exporting PyTorch models to ONNX format `__ guide to learn how to export models from PyTorch to ONNX. -2. Follow :doc:`Convert an ONNX model ` chapter to produce OpenVINO model. +2. Follow the :doc:`Convert an ONNX model ` guide to produce OpenVINO IR. Here is an illustration of using these two steps together: diff --git a/docs/articles_en/openvino-workflow/torch-compile.rst b/docs/articles_en/openvino-workflow/torch-compile.rst index 8c6016bfd4742f..d398704a819edc 100644 --- a/docs/articles_en/openvino-workflow/torch-compile.rst +++ b/docs/articles_en/openvino-workflow/torch-compile.rst @@ -5,7 +5,8 @@ PyTorch Deployment via "torch.compile" The ``torch.compile`` feature enables you to use OpenVINO for PyTorch-native applications. It speeds up PyTorch code by JIT-compiling it into optimized kernels. -By default, Torch code runs in eager-mode, but with the use of ``torch.compile`` it goes through the following steps: +By default, Torch code runs in eager-mode, but with the use of ``torch.compile`` it goes +through the following steps: 1. **Graph acquisition** - the model is rewritten as blocks of subgraphs that are either: diff --git a/docs/dev/pypi_publish/pypi-openvino-rt.md b/docs/dev/pypi_publish/pypi-openvino-rt.md index 854984ed2a0734..642eb12d65e8f9 100644 --- a/docs/dev/pypi_publish/pypi-openvino-rt.md +++ b/docs/dev/pypi_publish/pypi-openvino-rt.md @@ -6,8 +6,8 @@ Intel® Distribution of OpenVINO™ toolkit is an open-source toolkit for optimizing and deploying AI inference. It can be used to develop applications and solutions based on deep learning tasks, such as: emulation of human vision, automatic speech recognition, natural language processing, -recommendation systems, etc. It provides high-performance and rich deployment options, from -edge to cloud. +recommendation systems, image generation, etc. It provides high-performance and rich deployment +options, from edge to cloud. If you have chosen a model, you can integrate it with your application through OpenVINO™ and deploy it on various devices. The OpenVINO™ Python package includes a set of libraries for easy @@ -26,7 +26,7 @@ versions. The complete list of supported hardware is available on the ## Install OpenVINO™ -### Step 1. Set Up Python Virtual Environment +### Step 1. Set up Python virtual environment Use a virtual environment to avoid dependency conflicts. To create a virtual environment, use the following commands: @@ -43,7 +43,7 @@ python3 -m venv openvino_env > **NOTE**: On Linux and macOS, you may need to [install pip](https://pip.pypa.io/en/stable/installation/). -### Step 2. Activate the Virtual Environment +### Step 2. Activate the virtual environment On Windows: ```sh @@ -55,24 +55,23 @@ On Linux and macOS: source openvino_env/bin/activate ``` -### Step 3. Set Up and Update PIP to the Highest Version +### Step 3. Set up PIP and update it to the highest version -Run the command below: +Run the command: ```sh python -m pip install --upgrade pip ``` -### Step 4. Install the Package +### Step 4. Install the package -Run the command below:
- - ```sh - pip install openvino - ``` +Run the command: +```sh +pip install openvino +``` -### Step 5. Verify that the Package Is Installed +### Step 5. Verify that the package is installed -Run the command below: +Run the command: ```sh python -c "from openvino import Core; print(Core().available_devices)" ``` @@ -88,22 +87,22 @@ If installation was successful, you will see the list of available devices. Description - OpenVINO Runtime + OpenVINO Runtime `openvino package` OpenVINO Runtime is a set of C++ libraries with C and Python bindings providing a common API to deliver inference solutions on the platform of your choice. Use the OpenVINO Runtime API to read PyTorch, TensorFlow, TensorFlow Lite, ONNX, and PaddlePaddle models and execute them on preferred devices. OpenVINO Runtime uses a plugin architecture and includes the following plugins: - CPU, - GPU, - Auto Batch, - Auto, - Hetero, + CPU, + GPU, + Auto Batch, + Auto, + Hetero, - OpenVINO Model Converter (OVC) + OpenVINO Model Converter (OVC) `ovc` OpenVINO Model Converter converts models that were trained in popular frameworks to a format usable by OpenVINO components.
Supported frameworks include ONNX, TensorFlow, @@ -111,7 +110,7 @@ If installation was successful, you will see the list of available devices. - Benchmark Tool + Benchmark Tool `benchmark_app` Benchmark Application** allows you to estimate deep learning inference performance on supported devices for synchronous and asynchronous modes. @@ -122,8 +121,8 @@ If installation was successful, you will see the list of available devices. ## Troubleshooting -For general troubleshooting steps and issues, see -[Troubleshooting Guide for OpenVINO Installation](https://docs.openvino.ai/2024/get-started/troubleshooting-install-config.html). +For general troubleshooting, see the +[Troubleshooting Guide for OpenVINO Installation](https://docs.openvino.ai/2025/get-started/troubleshooting-install-config.html). The following sections also provide explanations to several error messages. ### Errors with Installing via PIP for Users in China @@ -145,11 +144,11 @@ the [C++ redistributable (.exe)](https://aka.ms/vs/17/release/vc_redist.x64.exe) You can also view a full download list on the [official support page](https://docs.microsoft.com/en-us/cpp/windows/latest-supported-vc-redist). -### ImportError: libpython3.8.so.1.0: cannot open shared object file: No such file or directory +### ImportError: libpython3.10.so.1.0: cannot open shared object file: No such file or directory To resolve missing external dependency on Ubuntu*, execute the following command: ```sh -sudo apt-get install libpython3.8 +sudo apt-get install libpython3.10 ``` ## Additional Resources @@ -159,7 +158,7 @@ sudo apt-get install libpython3.8 - [OpenVINO™ Notebooks](https://github.com/openvinotoolkit/openvino_notebooks) - [OpenVINO Installation Selector Tool](https://www.intel.com/content/www/us/en/developer/tools/openvino-toolkit/download.html) -Copyright © 2018-2024 Intel Corporation +Copyright © 2018-2025 Intel Corporation > **LEGAL NOTICE**: Your use of this software and any required dependent software (the “Software Package”) is subject to the terms and conditions of the [Apache 2.0 License](https://www.apache.org/licenses/LICENSE-2.0.html) for the Software Package, diff --git a/docs/sphinx_setup/_static/conformance_files/pytorch_ops.csv b/docs/sphinx_setup/_static/conformance_files/pytorch_ops.csv index a5836fb9b1f946..93910ec63cf4a2 100644 --- a/docs/sphinx_setup/_static/conformance_files/pytorch_ops.csv +++ b/docs/sphinx_setup/_static/conformance_files/pytorch_ops.csv @@ -1,3 +1,4 @@ +Operation Name,Limitation aten::__and__, aten::__derive_index, aten::__getitem__, @@ -39,7 +40,7 @@ aten::amax, aten::amin, aten::aminmax, aten::any, -aten::append , Supported in limited set of patterns +aten::append,Supported in limited set of patterns aten::arange, aten::argmax, aten::argmin, @@ -65,7 +66,7 @@ aten::bitwise_or, aten::bitwise_xor, aten::bmm, aten::Bool, -aten::broadcast_tensors , Supported in limited set of patterns +aten::broadcast_tensors,Supported in limited set of patterns aten::broadcast_to, aten::bucketize, aten::cat, @@ -75,7 +76,7 @@ aten::ceil_, aten::celu, aten::celu_, aten::channel_shuffle, -aten::chunk , Supported in limited set of patterns +aten::chunk,Supported in limited set of patterns aten::clamp, aten::clamp_, aten::clamp_max, @@ -86,7 +87,7 @@ aten::clip, aten::clip_, aten::clone, aten::col2im, -aten::complex , Supported in limited set of patterns +aten::complex,Supported in limited set of patterns aten::concat, aten::contiguous, aten::conv1d, @@ -112,7 +113,7 @@ aten::div_, aten::dot, aten::dropout, aten::dropout_, -aten::einsum , Supported in limited set of patterns +aten::einsum,Supported in limited set of patterns aten::elu, aten::elu_, aten::embedding, @@ -134,8 +135,8 @@ aten::eye, aten::fake_quantize_per_channel_affine, aten::fake_quantize_per_tensor_affine, aten::feature_dropout, -aten::fft_irfftn , Supported in limited set of patterns -aten::fft_rfftn , Supported in limited set of patterns +aten::fft_irfftn,Supported in limited set of patterns +aten::fft_rfftn,Supported in limited set of patterns aten::fill, aten::fill_, aten::fill_diagonal_, @@ -167,8 +168,8 @@ aten::hardswish_, aten::hardtanh, aten::hardtanh_, aten::im2col, -aten::imag , Supported in limited set of patterns -aten::index , Supported in limited set of patterns +aten::imag,Supported in limited set of patterns +aten::index,Supported in limited set of patterns aten::index_add, aten::index_add_, aten::index_copy_, @@ -250,7 +251,7 @@ aten::new_full, aten::new_ones, aten::new_zeros, aten::nonzero, -aten::nonzero_numpy , Supported in limited set of patterns +aten::nonzero_numpy,Supported in limited set of patterns aten::norm, aten::normal, aten::normal_, @@ -276,10 +277,10 @@ aten::rand_like, aten::randint, aten::randn, aten::randn_like, -aten::real , Supported in limited set of patterns +aten::real,Supported in limited set of patterns aten::reciprocal, aten::reciprocal_, -aten::reflection_pad2d , Supported in limited set of patterns +aten::reflection_pad2d,Supported in limited set of patterns aten::relu, aten::relu6, aten::relu6_, @@ -323,13 +324,13 @@ aten::slice, aten::softmax, aten::softplus, aten::sort, -aten::split , Supported in limited set of patterns -aten::split_with_sizes , Supported in limited set of patterns +aten::split,Supported in limited set of patterns +aten::split_with_sizes,Supported in limited set of patterns aten::sqrt, aten::sqrt_, aten::square, aten::squeeze, -aten::stack , Supported in limited set of patterns +aten::stack,Supported in limited set of patterns aten::std, aten::std_mean, aten::sub, @@ -344,7 +345,7 @@ aten::tan_, aten::tanh, aten::tanh_, aten::tensor, -aten::tensor_split , Supported in limited set of patterns +aten::tensor_split,Supported in limited set of patterns aten::tile, aten::to, aten::topk, @@ -354,7 +355,7 @@ aten::tril_, aten::triu, aten::triu_, aten::type_as, -aten::unbind , Supported in limited set of patterns +aten::unbind,Supported in limited set of patterns aten::unflatten, aten::unfold, aten::unsqueeze, @@ -379,21 +380,21 @@ ov_ext::embedding, ov_ext::linear, prim::Constant, prim::device, -prim::DictConstruct , Supported in limited set of patterns +prim::DictConstruct,Supported in limited set of patterns prim::GetAttr, prim::If, prim::is_cuda, prim::ListConstruct, prim::ListUnpack, prim::Loop, -prim::max , Supported in limited set of patterns -prim::min , Supported in limited set of patterns +prim::max,Supported in limited set of patterns +prim::min,Supported in limited set of patterns prim::NumToTensor, prim::PythonOp, prim::requires_grad, -prim::TupleConstruct , Supported in limited set of patterns +prim::TupleConstruct,Supported in limited set of patterns prim::TupleIndex, -prim::TupleUnpack , Supported in limited set of patterns +prim::TupleUnpack,Supported in limited set of patterns prim::type, quantized::add, quantized::add_relu, diff --git a/docs/sphinx_setup/_static/download/GenAI_Quick_Start_Guide.pdf b/docs/sphinx_setup/_static/download/GenAI_Quick_Start_Guide.pdf index 13edfc8f0b7bc2..5f24d28643598e 100644 Binary files a/docs/sphinx_setup/_static/download/GenAI_Quick_Start_Guide.pdf and b/docs/sphinx_setup/_static/download/GenAI_Quick_Start_Guide.pdf differ diff --git a/docs/sphinx_setup/_static/html/footer.html b/docs/sphinx_setup/_static/html/footer.html index d2ed2756109938..75e046fa72fefc 100644 --- a/docs/sphinx_setup/_static/html/footer.html +++ b/docs/sphinx_setup/_static/html/footer.html @@ -106,7 +106,7 @@
    - ©2024 Intel Corporation + ©2025 Intel Corporation
    OpenVINO diff --git a/docs/sphinx_setup/conf.py b/docs/sphinx_setup/conf.py index 8bf8438fb5e2c2..ffe00effa610b4 100644 --- a/docs/sphinx_setup/conf.py +++ b/docs/sphinx_setup/conf.py @@ -187,13 +187,11 @@ html_static_path = ['_static'] html_css_files = [ - 'css/custom.css', 'css/openvino_sphinx_theme.css', 'css/button.css', 'css/input.css', 'css/textfield.css', 'css/tabs.css', - 'css/coveo_custom.css', 'https://cdn.jsdelivr.net/npm/@splidejs/splide@4.1.4/dist/css/splide.min.css', ] diff --git a/src/core/include/openvino/pass/pattern/op/pattern.hpp b/src/core/include/openvino/pass/pattern/op/pattern.hpp index e9bb427d906b43..08afddd43a1c09 100644 --- a/src/core/include/openvino/pass/pattern/op/pattern.hpp +++ b/src/core/include/openvino/pass/pattern/op/pattern.hpp @@ -104,4 +104,24 @@ class OPENVINO_API Pattern : public Node { } // namespace op } // namespace pattern } // namespace pass +OPENVINO_API pass::pattern::op::ValuePredicate operator||(const std::function)>& a, + const std::function)>& b); +OPENVINO_API pass::pattern::op::ValuePredicate operator||(const pass::pattern::op::ValuePredicate& a, + const pass::pattern::op::ValuePredicate& b); +OPENVINO_API pass::pattern::op::ValuePredicate operator||(const pass::pattern::op::NodePredicate& a, + const pass::pattern::op::NodePredicate& b); +OPENVINO_API pass::pattern::op::ValuePredicate operator||(const pass::pattern::op::ValuePredicate& a, + const pass::pattern::op::NodePredicate& b); +OPENVINO_API pass::pattern::op::ValuePredicate operator||(const pass::pattern::op::NodePredicate& a, + const pass::pattern::op::ValuePredicate& b); +OPENVINO_API pass::pattern::op::ValuePredicate operator&&(const std::function)>& a, + const std::function)>& b); +OPENVINO_API pass::pattern::op::ValuePredicate operator&&(const pass::pattern::op::ValuePredicate& a, + const pass::pattern::op::ValuePredicate& b); +OPENVINO_API pass::pattern::op::ValuePredicate operator&&(const pass::pattern::op::NodePredicate& a, + const pass::pattern::op::NodePredicate& b); +OPENVINO_API pass::pattern::op::ValuePredicate operator&&(const pass::pattern::op::ValuePredicate& a, + const pass::pattern::op::NodePredicate& b); +OPENVINO_API pass::pattern::op::ValuePredicate operator&&(const pass::pattern::op::NodePredicate& a, + const pass::pattern::op::ValuePredicate& b); } // namespace ov diff --git a/src/core/src/pattern/op/pattern.cpp b/src/core/src/pattern/op/pattern.cpp index 4a1c7f367864da..826c6ff532ffb6 100644 --- a/src/core/src/pattern/op/pattern.cpp +++ b/src/core/src/pattern/op/pattern.cpp @@ -166,4 +166,61 @@ std::function)> all_of(const std::vector& value) { + return a(value) || b(value); + }; +} + +pass::pattern::op::ValuePredicate operator||(const std::function)>& a, + const std::function)>& b) { + return [=](const Output& value) { + return a(value) || b(value); + }; +} + +pass::pattern::op::ValuePredicate operator||(const pass::pattern::op::NodePredicate& a, + const pass::pattern::op::NodePredicate& b) { + return pass::pattern::op::as_value_predicate(a) || pass::pattern::op::as_value_predicate(b); +} + +pass::pattern::op::ValuePredicate operator||(const pass::pattern::op::ValuePredicate& a, + const pass::pattern::op::NodePredicate& b) { + return a || pass::pattern::op::as_value_predicate(b); +} + +pass::pattern::op::ValuePredicate operator||(const pass::pattern::op::NodePredicate& a, + const pass::pattern::op::ValuePredicate& b) { + return pass::pattern::op::as_value_predicate(a) || b; +} + +pass::pattern::op::ValuePredicate operator&&(const pass::pattern::op::ValuePredicate& a, + const pass::pattern::op::ValuePredicate& b) { + return [=](const Output& value) { + return a(value) && b(value); + }; +} + +pass::pattern::op::ValuePredicate operator&&(const std::function)>& a, + const std::function)>& b) { + return [=](const Output& value) { + return a(value) && b(value); + }; +} + +pass::pattern::op::ValuePredicate operator&&(const pass::pattern::op::NodePredicate& a, + const pass::pattern::op::NodePredicate& b) { + return pass::pattern::op::as_value_predicate(a) && pass::pattern::op::as_value_predicate(b); +} + +pass::pattern::op::ValuePredicate operator&&(const pass::pattern::op::ValuePredicate& a, + const pass::pattern::op::NodePredicate& b) { + return a && pass::pattern::op::as_value_predicate(b); +} + +pass::pattern::op::ValuePredicate operator&&(const pass::pattern::op::NodePredicate& a, + const pass::pattern::op::ValuePredicate& b) { + return pass::pattern::op::as_value_predicate(a) && b; +} } // namespace ov diff --git a/src/core/tests/pattern.cpp b/src/core/tests/pattern.cpp index 5bb961e57db1c2..39d30c3fd27492 100644 --- a/src/core/tests/pattern.cpp +++ b/src/core/tests/pattern.cpp @@ -1258,3 +1258,35 @@ TEST(pattern, pattern_optional_root) { // Should perfectly match ASSERT_TRUE(tm.match(pattern_relu, model_relu)); } + +TEST(pattern, pattern_predicate_operator) { + // Sample model + PartialShape shape{2, 2}; + auto model_param1 = std::make_shared(element::i32, shape); + auto model_param2 = std::make_shared(element::i32, shape); + auto model_add = std::make_shared(model_param1->output(0), model_param2->output(0)); + auto model_result = std::make_shared(model_add->output(0)); + + TestMatcher tm; + + ASSERT_TRUE(tm.match(ov::pass::pattern::any_input(ov::pass::pattern::rank_equals(2) || + ov::pass::pattern::type_matches(element::Type_t::boolean)), + model_add)); + ASSERT_TRUE(tm.match(ov::pass::pattern::any_input(ov::pass::pattern::type_matches(element::Type_t::boolean) || + ov::pass::pattern::rank_equals(2)), + model_add)); + + ASSERT_FALSE(tm.match(ov::pass::pattern::any_input(ov::pass::pattern::rank_equals(2) && + ov::pass::pattern::type_matches(element::Type_t::boolean)), + model_add)); + ASSERT_FALSE(tm.match(ov::pass::pattern::any_input(ov::pass::pattern::type_matches(element::Type_t::boolean) && + ov::pass::pattern::rank_equals(2)), + model_add)); + + ASSERT_TRUE(tm.match(ov::pass::pattern::any_input(ov::pass::pattern::rank_equals(2) && + ov::pass::pattern::type_matches(element::Type_t::i32)), + model_add)); + ASSERT_TRUE(tm.match(ov::pass::pattern::any_input(ov::pass::pattern::type_matches(element::Type_t::i32) && + ov::pass::pattern::rank_equals(2)), + model_add)); +} diff --git a/src/frontends/jax/src/op_table.cpp b/src/frontends/jax/src/op_table.cpp index 9c492dfa3e119d..6ae0e6adc7c469 100644 --- a/src/frontends/jax/src/op_table.cpp +++ b/src/frontends/jax/src/op_table.cpp @@ -19,6 +19,7 @@ #include "openvino/op/not_equal.hpp" #include "openvino/op/reduce_max.hpp" #include "openvino/op/reduce_sum.hpp" +#include "openvino/op/sigmoid.hpp" #include "openvino/op/sqrt.hpp" #include "openvino/op/subtract.hpp" #include "openvino/op/tanh.hpp" @@ -94,6 +95,7 @@ const std::map get_supported_ops_jaxpr() { {"rsqrt", op::translate_rsqrt}, {"reshape", op::translate_reshape}, {"select_n", op::translate_select_n}, + {"logistic", op::translate_1to1_match_1_input}, {"slice", op::translate_slice}, {"square", op::translate_square}, {"sqrt", op::translate_1to1_match_1_input}, diff --git a/src/inference/src/os/lin/lin_system_conf.cpp b/src/inference/src/os/lin/lin_system_conf.cpp index 29c8bfddbd1ca4..583c194c720c84 100644 --- a/src/inference/src/os/lin/lin_system_conf.cpp +++ b/src/inference/src/os/lin/lin_system_conf.cpp @@ -221,14 +221,18 @@ CPU::CPU() { } else { _processors = valid_cpu_mapping_table.size(); _cpu_mapping_table.swap(valid_cpu_mapping_table); + int cur_numa_nodes = _numa_nodes; + int cur_cores = _cores; { std::lock_guard lock{_cpu_mutex}; update_valid_processor_linux(std::move(phy_core_list), - _numa_nodes, - _cores, + cur_numa_nodes, + cur_cores, _proc_type_table, _cpu_mapping_table); } + _cores = cur_cores; + _numa_nodes = cur_numa_nodes; return 0; } }; diff --git a/src/plugins/intel_cpu/CMakeLists.txt b/src/plugins/intel_cpu/CMakeLists.txt index aa6ce49a051e00..c88ec8fa2a1331 100644 --- a/src/plugins/intel_cpu/CMakeLists.txt +++ b/src/plugins/intel_cpu/CMakeLists.txt @@ -283,16 +283,16 @@ target_include_directories(${TARGET_NAME} PRIVATE $> jit_prelu_emitter::get_supported_precisions( const std::shared_ptr& node) { - return {{element::f32}}; + return {{element::f32, element::f32}}; } void jit_prelu_emitter::emit_impl(const std::vector& in_vec_idxs, diff --git a/src/plugins/intel_cpu/src/emitters/plugin/x64/jit_dnnl_emitters.cpp b/src/plugins/intel_cpu/src/emitters/plugin/x64/jit_dnnl_emitters.cpp index 51e801208b927c..76ee7f9740c8c1 100644 --- a/src/plugins/intel_cpu/src/emitters/plugin/x64/jit_dnnl_emitters.cpp +++ b/src/plugins/intel_cpu/src/emitters/plugin/x64/jit_dnnl_emitters.cpp @@ -46,13 +46,18 @@ jit_dnnl_emitter::jit_dnnl_emitter(jit_generator* host, void jit_dnnl_emitter::set_injector() { if (host_isa_ == cpu::x64::sse41) { eltwise_injector_sse42 = - std::make_shared>(h, kind, alpha, beta, 1.f); + std::make_shared>(h, kind, alpha, beta, 1.f, data_type::f32); } else if (host_isa_ == cpu::x64::avx2) { eltwise_injector_avx2 = - std::make_shared>(h, kind, alpha, beta, 1.f); + std::make_shared>(h, kind, alpha, beta, 1.f, data_type::f32); } else if (host_isa_ == cpu::x64::avx512_core) { eltwise_injector_avx512_core = - std::make_shared>(h, kind, alpha, beta, 1.f); + std::make_shared>(h, + kind, + alpha, + beta, + 1.f, + data_type::f32); } else { OV_CPU_JIT_EMITTER_THROW("Unsupported ISA ", host_isa_); } diff --git a/src/plugins/intel_cpu/src/emitters/plugin/x64/jit_dnnl_emitters.hpp b/src/plugins/intel_cpu/src/emitters/plugin/x64/jit_dnnl_emitters.hpp index 22e003ad261555..3b569308c05a26 100644 --- a/src/plugins/intel_cpu/src/emitters/plugin/x64/jit_dnnl_emitters.hpp +++ b/src/plugins/intel_cpu/src/emitters/plugin/x64/jit_dnnl_emitters.hpp @@ -42,11 +42,9 @@ class jit_dnnl_emitter : public jit_emitter { float alpha{0.f}; float beta{0.f}; - std::shared_ptr> - eltwise_injector_sse42; - std::shared_ptr> - eltwise_injector_avx2; - std::shared_ptr> + std::shared_ptr> eltwise_injector_sse42; + std::shared_ptr> eltwise_injector_avx2; + std::shared_ptr> eltwise_injector_avx512_core; private: diff --git a/src/plugins/intel_cpu/src/emitters/snippets/aarch64/cpu_generator.cpp b/src/plugins/intel_cpu/src/emitters/snippets/aarch64/cpu_generator.cpp index 0f6b2c24c13df7..cf89181e2a7979 100644 --- a/src/plugins/intel_cpu/src/emitters/snippets/aarch64/cpu_generator.cpp +++ b/src/plugins/intel_cpu/src/emitters/snippets/aarch64/cpu_generator.cpp @@ -13,7 +13,13 @@ #include "emitters/snippets/cpu_runtime_configurator.hpp" #include "emitters/utils.hpp" #include "jit_snippets_emitters.hpp" +#include "openvino/core/type.hpp" +#include "openvino/op/prelu.hpp" +#include "openvino/op/round.hpp" +#include "openvino/op/sqrt.hpp" #include "openvino/opsets/opset13.hpp" +#include "snippets/emitter.hpp" +#include "snippets/lowered/expression.hpp" #include "snippets/snippets_isa.hpp" #include "transformations/cpu_opset/common/op/swish_cpu.hpp" #include "transformations/snippets/common/op/fused_mul_add.hpp" @@ -44,7 +50,7 @@ namespace ov { { \ [this](const snippets::lowered::ExpressionPtr& expr) -> std::shared_ptr { \ const auto& n = expr->get_node(); \ - const auto& gelu = std::dynamic_pointer_cast(n); \ + const auto& gelu = ov::as_type_ptr(n); \ if (gelu == nullptr) { \ OPENVINO_THROW("Can't cast to ov::op::v7::Gelu"); \ } \ @@ -73,6 +79,37 @@ namespace ov { } \ } +#define CREATE_ROUND_V5_EMITTER(e_type_from_zero, e_type_even) \ + { \ + [this](const snippets::lowered::ExpressionPtr& expr) -> std::shared_ptr { \ + const auto& n = expr->get_node(); \ + const auto& round = ov::as_type_ptr(n); \ + if (round == nullptr) { \ + OPENVINO_THROW("Can't cast to ov::op::v5::Round"); \ + } \ + const auto roundingMode = round->get_mode(); \ + if (roundingMode == ov::op::v5::Round::RoundMode::HALF_AWAY_FROM_ZERO) { \ + return std::make_shared(h.get(), isa, n); \ + } else if (roundingMode == ov::op::v5::Round::RoundMode::HALF_TO_EVEN) { \ + return std::make_shared(h.get(), isa, n); \ + } else { \ + OPENVINO_THROW("Unsupported Round mode"); \ + } \ + }, \ + [](const std::shared_ptr& n) -> std::set> { \ + const auto& round = std::dynamic_pointer_cast(n); \ + if (round == nullptr) { \ + OPENVINO_THROW("Can't cast to ov::op::v5::Round"); \ + } \ + if (round->get_mode() == ov::op::v5::Round::RoundMode::HALF_AWAY_FROM_ZERO) { \ + return e_type_from_zero::get_supported_precisions(n); \ + } else if (round->get_mode() == ov::op::v5::Round::RoundMode::HALF_TO_EVEN) { \ + return e_type_even::get_supported_precisions(n); \ + } \ + OPENVINO_THROW("Unsupported Round mode"); \ + } \ + } + class jit_snippet : public dnnl::impl::cpu::aarch64::jit_generator { public: DECLARE_CPU_JIT_AUX_FUNCTIONS(jit_snippet) @@ -136,6 +173,12 @@ CPUTargetMachine::CPUTargetMachine(dnnl::impl::cpu::aarch64::cpu_isa_t host_isa) jitters[op::v1::Multiply::get_type_info_static()] = CREATE_CPU_EMITTER(jit_multiply_emitter); jitters[op::v1::Subtract::get_type_info_static()] = CREATE_CPU_EMITTER(jit_subtract_emitter); + // Comparison ops + jitters[op::v1::Equal::get_type_info_static()] = CREATE_CPU_EMITTER(jit_equal_emitter); + jitters[op::v1::Greater::get_type_info_static()] = CREATE_CPU_EMITTER(jit_greater_emitter); + jitters[op::v1::GreaterEqual::get_type_info_static()] = CREATE_CPU_EMITTER(jit_greater_equal_emitter); + jitters[op::v1::LessEqual::get_type_info_static()] = CREATE_CPU_EMITTER(jit_less_equal_emitter); + // unary jitters[ov::op::v0::Abs::get_type_info_static()] = CREATE_CPU_EMITTER(jit_abs_emitter); jitters[ov::op::v0::Ceiling::get_type_info_static()] = CREATE_CPU_EMITTER(jit_ceiling_emitter); @@ -149,8 +192,12 @@ CPUTargetMachine::CPUTargetMachine(dnnl::impl::cpu::aarch64::cpu_isa_t host_isa) CREATE_GELU_V7_EMITTER(jit_gelu_erf_emitter, jit_gelu_tanh_emitter); jitters[ov::op::v4::HSwish::get_type_info_static()] = CREATE_CPU_EMITTER(jit_hswish_emitter); jitters[ov::op::v4::Mish::get_type_info_static()] = CREATE_CPU_EMITTER(jit_mish_emitter); + jitters[ov::op::v0::PRelu::get_type_info_static()] = CREATE_CPU_EMITTER(jit_prelu_emitter); jitters[ov::op::v0::Relu::get_type_info_static()] = CREATE_CPU_EMITTER(jit_relu_emitter); + jitters[ov::op::v5::Round::get_type_info_static()] = + CREATE_ROUND_V5_EMITTER(jit_round_half_away_from_zero_emitter, jit_round_half_to_even_emitter); jitters[ov::op::v0::Sigmoid::get_type_info_static()] = CREATE_CPU_EMITTER(jit_sigmoid_emitter); + jitters[ov::op::v0::Sqrt::get_type_info_static()] = CREATE_CPU_EMITTER(jit_sqrt_emitter); jitters[ov::intel_cpu::SwishNode::get_type_info_static()] = CREATE_CPU_EMITTER(jit_swish_emitter); jitters[ov::op::v0::Tanh::get_type_info_static()] = CREATE_CPU_EMITTER(jit_tanh_emitter); diff --git a/src/plugins/intel_cpu/src/nodes/adaptive_pooling.cpp b/src/plugins/intel_cpu/src/nodes/adaptive_pooling.cpp index f4af11b0f2362a..c0627ff43bf6ba 100644 --- a/src/plugins/intel_cpu/src/nodes/adaptive_pooling.cpp +++ b/src/plugins/intel_cpu/src/nodes/adaptive_pooling.cpp @@ -191,7 +191,7 @@ void AdaptivePooling::execute(dnnl::stream strm) { auto dstStrides = getChildEdgeAt(0)->getMemory().getDescWithType()->getStrides(); // unified strides array - const size_t tailDimsOffset = (isTailCFmt ? -1 : 0); + const ptrdiff_t tailDimsOffset = (isTailCFmt ? -1 : 0); const size_t inStrides[5] = {srcStrides[0], (isTailCFmt ? 1 : srcStrides[1]), (spatialDimsCount == 3 ? srcStrides[2 + tailDimsOffset] : 0), diff --git a/src/plugins/intel_cpu/src/nodes/batch_to_space.cpp b/src/plugins/intel_cpu/src/nodes/batch_to_space.cpp index 50665c083ec930..098e93eeb19a60 100644 --- a/src/plugins/intel_cpu/src/nodes/batch_to_space.cpp +++ b/src/plugins/intel_cpu/src/nodes/batch_to_space.cpp @@ -156,6 +156,7 @@ void BatchToSpace::batchToSpaceKernel() { size_t channels = (inShape5D[1] / blockSize); channels = channels == 0 ? 1 : channels; const size_t workAmount = inShape5D[0] * channels; + OPENVINO_ASSERT(workAmount > 0, errorPrefix, " has unsupported work amount == 0"); parallel_nt(0, [&](const int ithr, const int nthr) { size_t start(0lu), end(0lu); diff --git a/src/plugins/intel_cpu/src/nodes/bin_conv.cpp b/src/plugins/intel_cpu/src/nodes/bin_conv.cpp index 336a370374a9f9..68952e6ccc5199 100644 --- a/src/plugins/intel_cpu/src/nodes/bin_conv.cpp +++ b/src/plugins/intel_cpu/src/nodes/bin_conv.cpp @@ -66,11 +66,12 @@ struct jit_uni_bin_conv_kernel_f32 : public jit_uni_bin_conv_kernel, public jit_ for (int i = 0; i < end_idx; i++) { auto& post_op = p.entry_[i]; if (post_op.is_eltwise()) { - eltwise_injectors.push_back(std::make_shared>(this, - post_op.eltwise, - true, - eltwise_reserved, - mask_post_op_reserved)); + eltwise_injectors.push_back(std::make_shared>(this, + post_op.eltwise, + data_type::f32, + true, + eltwise_reserved, + mask_post_op_reserved)); } else if (post_op.is_depthwise()) { depthwise_injectors.push_back( std::make_shared>(this, post_op, mask_post_op_reserved)); @@ -217,7 +218,7 @@ struct jit_uni_bin_conv_kernel_f32 : public jit_uni_bin_conv_kernel, public jit_ Xbyak::Label l_table; - nstl::vector>> eltwise_injectors; + nstl::vector>> eltwise_injectors; nstl::vector>> depthwise_injectors; void cvt2ps(dnnl::memory::data_type type_in, Vmm vmm_in, const Xbyak::Operand& op, bool scalar_load) { diff --git a/src/plugins/intel_cpu/src/nodes/common/softmax.cpp b/src/plugins/intel_cpu/src/nodes/common/softmax.cpp index 0fcc87f8978752..888bc40412401b 100644 --- a/src/plugins/intel_cpu/src/nodes/common/softmax.cpp +++ b/src/plugins/intel_cpu/src/nodes/common/softmax.cpp @@ -16,6 +16,7 @@ #include "utils/bfloat16.hpp" using namespace dnnl; +using namespace dnnl::impl; using namespace dnnl::impl::cpu; using namespace dnnl::impl::cpu::x64; using namespace dnnl::impl::utils; @@ -68,7 +69,7 @@ struct jit_uni_softmax_kernel_f32 : public jit_uni_softmax_kernel, public jit_ge void generate() override { exp_injector.reset( - new jit_uni_eltwise_injector_f32(this, dnnl::impl::alg_kind::eltwise_exp, 0.f, 0.f, 1.0f)); + new jit_uni_eltwise_injector(this, dnnl::impl::alg_kind::eltwise_exp, 0.f, 0.f, 1.0f, data_type::f32)); if (mayiuse(avx512_core)) uni_vcvtneps2bf16.reset(new jit_uni_vcvtneps2bf16(this, isa)); @@ -200,7 +201,7 @@ struct jit_uni_softmax_kernel_f32 : public jit_uni_softmax_kernel, public jit_ge std::unique_ptr uni_vcvtneps2bf16; - std::shared_ptr> exp_injector; + std::shared_ptr> exp_injector; jit_softmax_config_params jcp_; diff --git a/src/plugins/intel_cpu/src/nodes/interpolate.cpp b/src/plugins/intel_cpu/src/nodes/interpolate.cpp index beb53cb89a831e..0cc3c71cf0319b 100644 --- a/src/plugins/intel_cpu/src/nodes/interpolate.cpp +++ b/src/plugins/intel_cpu/src/nodes/interpolate.cpp @@ -73,11 +73,12 @@ struct jit_uni_interpolate_kernel_f32 : public jit_uni_interpolate_kernel, publi for (int i = 0; i < p.len(); i++) { auto& post_op = p.entry_[i]; if (post_op.is_eltwise()) { - eltwise_injectors.push_back(std::make_shared>(this, - post_op.eltwise.alg, - post_op.eltwise.alpha, - post_op.eltwise.beta, - 1.f)); + eltwise_injectors.push_back(std::make_shared>(this, + post_op.eltwise.alg, + post_op.eltwise.alpha, + post_op.eltwise.beta, + 1.f, + data_type::f32)); } else if (post_op.is_depthwise()) { depthwise_injectors.push_back(std::make_shared>(this, post_op)); } else if (post_op.is_quantization()) { @@ -275,7 +276,7 @@ struct jit_uni_interpolate_kernel_f32 : public jit_uni_interpolate_kernel, publi std::vector store_pool_vec_idxs; std::vector load_pool_gpr_idxs; - std::vector>> eltwise_injectors; + std::vector>> eltwise_injectors; std::vector>> depthwise_injectors; std::vector>> quantization_injectors; diff --git a/src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/common.hpp b/src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/common.hpp index 63cbbb4464ee92..cb1cd24f840bfd 100644 --- a/src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/common.hpp +++ b/src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/common.hpp @@ -39,7 +39,14 @@ static constexpr size_t vec_len_f32_neon = vec_len_neon / sizeof(float); static constexpr size_t vec_len_f16_neon = vec_len_neon / sizeof(ov::float16); #if defined(HAVE_SVE) -static constexpr size_t vec_len_f32_sve = svcntw(); +inline size_t vec_len_f32_sve() { + static size_t len = svcntw(); + return len; +} +inline size_t vec_len_f16_sve() { + static size_t len = svcnth(); + return len; +} #endif #ifdef HAVE_AVX512F @@ -403,6 +410,28 @@ inline void __vst1q_f32(ov::bfloat16* a, float32x4_t b) { #endif #if defined(__ARM_FEATURE_FP16_VECTOR_ARITHMETIC) +# if defined(HAVE_SVE) +inline svfloat16_t exp_ps_sve_f16(svbool_t& pg, svfloat16_t& src) { + svbool_t pg_f32 = svtrn1_b16(pg, svpfalse()); + + // Extract lower and upper halves of src into two separate vecs and convert + svfloat16_t zero = svdup_n_f16(0.0); + svfloat16_t low_f16 = svtrn1_f16(src, zero); + svfloat16_t high_f16 = svtrn2_f16(src, zero); + svfloat32_t low_f32 = svcvt_f32_f16_z(pg, low_f16); + svfloat32_t high_f32 = svcvt_f32_f16_z(pg, high_f16); + + // Perform exp and convert back to f16 + svfloat32_t low_exp_f32 = exp_ps_sve(pg_f32, low_f32); + svfloat32_t high_exp_f32 = exp_ps_sve(pg_f32, high_f32); + svfloat16_t low_exp_f16 = svcvt_f16_f32_z(pg_f32, low_exp_f32); + svfloat16_t high_exp_f16 = svcvt_f16_f32_z(pg_f32, high_exp_f32); + + // Interleave both to get final result + svfloat16_t res = svtrn1_f16(low_exp_f16, high_exp_f16); + return res; +} +# else inline float16x8_t exp_ps_neon_f16(float16x8_t x) { const float32x4_t x_high = vcvt_f32_f16(vget_high_f16(x)); const float32x4_t x_low = vcvt_f32_f16(vget_low_f16(x)); @@ -411,6 +440,7 @@ inline float16x8_t exp_ps_neon_f16(float16x8_t x) { const float16x8_t res = vcombine_f16(vcvt_f16_f32(exp_ps_neon_f32(x_low)), vcvt_f16_f32(exp_ps_neon_f32(x_high))); return res; } +# endif inline float16_t hsum(float16x8_t vec) { float16x4_t sum1 = vpadd_f16(vget_low_f16(vec), vget_high_f16(vec)); float16x4_t sum2 = vpadd_f16(sum1, sum1); diff --git a/src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/mha_single_token.cpp b/src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/mha_single_token.cpp index 5a6f0d66f1f221..8f56da752b6f5c 100644 --- a/src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/mha_single_token.cpp +++ b/src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/mha_single_token.cpp @@ -63,13 +63,36 @@ void cvt_copy(TA* dst, TB* src, size_t n) { mm256_uni_storeu_ps(dst + i, vb); } #elif defined(OPENVINO_ARCH_ARM64) -# if defined(HAVE_SVE) +# if defined(__ARM_FEATURE_FP16_VECTOR_ARITHMETIC) + if (std::is_same::value && std::is_same::value) { +# if defined(HAVE_SVE) + size_t inc = vec_len_f16_sve(); + svbool_t pg = svptrue_b16(); + + while (i < n) { + if (n - i < vec_len_f16_sve()) { + inc = n - i; + pg = svwhilelt_b16(0, static_cast(inc)); + } + svfloat16_t b1 = svld1_f16(pg, reinterpret_cast(src + i)); + svst1_f16(pg, reinterpret_cast(dst + i), b1); + i += inc; + } +# else + for (; i + vec_len_f16_neon <= n; i += vec_len_f16_neon) { + auto vb1 = vld1q_f16(reinterpret_cast(src + i)); + vst1q_f16(reinterpret_cast(dst + i), vb1); + } +# endif + } +# else +# if defined(HAVE_SVE) auto _dst = reinterpret_cast(dst); - size_t inc = vec_len_f32_sve; + size_t inc = vec_len_f32_sve(); svbool_t pg = svptrue_b32(); while (i < n) { - if (n - i < vec_len_f32_sve) { + if (n - i < vec_len_f32_sve()) { inc = n - i; pg = svwhilelt_b32(0, static_cast(inc)); } @@ -77,20 +100,13 @@ void cvt_copy(TA* dst, TB* src, size_t n) { svst1_f32(pg, _dst + i, b1); i += inc; } -# else +# else if (std::is_same::value && std::is_same::value) { for (; i + vec_len_f32_neon <= n; i += vec_len_f32_neon) { float32x4_t vb1 = __vld1q_f32(src + i); __vst1q_f32(dst + i, vb1); } } -# if defined(__ARM_FEATURE_FP16_VECTOR_ARITHMETIC) - if (std::is_same::value && std::is_same::value) { - for (; i + vec_len_f16_neon <= n; i += vec_len_f16_neon) { - auto vb1 = vld1q_f16(reinterpret_cast(src + i)); - vst1q_f16(reinterpret_cast(dst + i), vb1); - } - } # endif # endif #endif @@ -122,11 +138,11 @@ static void attn_acc_value(float* out, float weight, T* v, size_t S, float* scal # if defined(HAVE_SVE) auto _v = reinterpret_cast(v); svfloat32_t attn_w_vec_fp32 = svdup_n_f32(weight); - size_t inc = vec_len_f32_sve; + size_t inc = vec_len_f32_sve(); svbool_t pg = svptrue_b32(); while (i < S) { - if (S - i < vec_len_f32_sve) { + if (S - i < vec_len_f32_sve()) { inc = S - i; pg = svwhilelt_b32(0, static_cast(inc)); } @@ -158,15 +174,35 @@ static void attn_acc_value(float* out, float weight, T* v, size_t S, float* scal template static void attn_acc_value(ov::float16* out, ov::float16 weight, T* v, size_t S, float* scale, float* zp) { size_t i = 0; - auto attn_w_vec_fp16 = vdupq_n_f16(weight); auto _v = reinterpret_cast(v); auto _out = reinterpret_cast(out); + +# if defined(HAVE_SVE) + svfloat16_t attn_w_vec_fp16 = svdup_n_f16(weight); + svbool_t pg = svptrue_b16(); + size_t inc = vec_len_f16_sve(); + + while (i < S) { + if (S - i < vec_len_f16_sve()) { + inc = S - i; + pg = svwhilelt_b16(0, static_cast(inc)); + } + svfloat16_t v_value = svld1_f16(pg, _v + i); + svfloat16_t v_out = svld1_f16(pg, _out + i); + + v_out = svmla_f16_m(pg, v_out, attn_w_vec_fp16, v_value); + svst1_f16(pg, _out + i, v_out); + i += inc; + } +# else + auto attn_w_vec_fp16 = vdupq_n_f16(weight); for (; i + vec_len_f16_neon <= S; i += vec_len_f16_neon) { auto v_value = vld1q_f16(_v + i); auto v_out = vld1q_f16(_out + i); v_out = vfmaq_f16(v_out, attn_w_vec_fp16, v_value); vst1q_f16(_out + i, v_out); } +# endif for (; i < S; i++) { out[i] += weight * v[i]; } @@ -404,30 +440,31 @@ static float sum_q_head(T* a, size_t n) { svfloat32_t sum2 = svdup_n_f32(0.0f); svfloat32_t sum3 = svdup_n_f32(0.0f); svbool_t pg = svptrue_b32(); + auto vec_len = vec_len_f32_sve(); - for (; i + 4 * vec_len_f32_sve <= n; i += 4 * vec_len_f32_sve) { + for (; i + 4 * vec_len <= n; i += 4 * vec_len) { svfloat32_t a0 = svld1_f32(pg, a + i); - svfloat32_t a1 = svld1_f32(pg, a + i + vec_len_f32_sve); - svfloat32_t a2 = svld1_f32(pg, a + i + vec_len_f32_sve * 2); - svfloat32_t a3 = svld1_f32(pg, a + i + vec_len_f32_sve * 3); + svfloat32_t a1 = svld1_f32(pg, a + i + vec_len); + svfloat32_t a2 = svld1_f32(pg, a + i + vec_len * 2); + svfloat32_t a3 = svld1_f32(pg, a + i + vec_len * 3); sum0 = svadd_f32_z(pg, a0, sum0); sum1 = svadd_f32_z(pg, a1, sum1); sum2 = svadd_f32_z(pg, a2, sum2); sum3 = svadd_f32_z(pg, a3, sum3); } - if (i + 2 * vec_len_f32_sve <= n) { + if (i + 2 * vec_len <= n) { svfloat32_t a0 = svld1_f32(pg, a + i); - svfloat32_t a1 = svld1_f32(pg, a + i + vec_len_f32_sve); + svfloat32_t a1 = svld1_f32(pg, a + i + vec_len); sum0 = svadd_f32_z(pg, a0, sum0); sum1 = svadd_f32_z(pg, a1, sum1); - i += 2 * vec_len_f32_sve; + i += 2 * vec_len; } - if (i + vec_len_f32_sve <= n) { + if (i + vec_len <= n) { svfloat32_t a0 = svld1_f32(pg, a + i); sum0 = svadd_f32_z(pg, a0, sum0); - i += vec_len_f32_sve; + i += vec_len; } // Process tail elements parallely as well (if any) if (i != n) { @@ -587,42 +624,43 @@ static float dot_product(TA* a, TB* b, size_t n, float* scale, float* zp, float* svfloat32_t sum1 = svdup_n_f32(0.0f); svfloat32_t sum2 = svdup_n_f32(0.0f); svfloat32_t sum3 = svdup_n_f32(0.0f); + auto vec_len = vec_len_f32_sve(); auto _a = reinterpret_cast(a); auto _b = reinterpret_cast(b); - for (; i + 4 * vec_len_f32_sve <= n; i += 4 * vec_len_f32_sve) { + for (; i + 4 * vec_len <= n; i += 4 * vec_len) { svfloat32_t a0 = svld1_f32(pg, _a + i); - svfloat32_t a1 = svld1_f32(pg, _a + i + vec_len_f32_sve); - svfloat32_t a2 = svld1_f32(pg, _a + i + vec_len_f32_sve * 2); - svfloat32_t a3 = svld1_f32(pg, _a + i + vec_len_f32_sve * 3); + svfloat32_t a1 = svld1_f32(pg, _a + i + vec_len); + svfloat32_t a2 = svld1_f32(pg, _a + i + vec_len * 2); + svfloat32_t a3 = svld1_f32(pg, _a + i + vec_len * 3); svfloat32_t b0 = svld1_f32(pg, _b + i); - svfloat32_t b1 = svld1_f32(pg, _b + i + vec_len_f32_sve); - svfloat32_t b2 = svld1_f32(pg, _b + i + vec_len_f32_sve * 2); - svfloat32_t b3 = svld1_f32(pg, _b + i + vec_len_f32_sve * 3); + svfloat32_t b1 = svld1_f32(pg, _b + i + vec_len); + svfloat32_t b2 = svld1_f32(pg, _b + i + vec_len * 2); + svfloat32_t b3 = svld1_f32(pg, _b + i + vec_len * 3); sum0 = svmla_f32_z(pg, sum0, a0, b0); sum1 = svmla_f32_z(pg, sum1, a1, b1); sum2 = svmla_f32_z(pg, sum2, a2, b2); sum3 = svmla_f32_z(pg, sum3, a3, b3); } - if (i + 2 * vec_len_f32_sve <= n) { + if (i + 2 * vec_len <= n) { svfloat32_t a0 = svld1_f32(pg, _a + i); - svfloat32_t a1 = svld1_f32(pg, _a + i + vec_len_f32_sve); + svfloat32_t a1 = svld1_f32(pg, _a + i + vec_len); svfloat32_t b0 = svld1_f32(pg, _b + i); - svfloat32_t b1 = svld1_f32(pg, _b + i + vec_len_f32_sve); + svfloat32_t b1 = svld1_f32(pg, _b + i + vec_len); sum0 = svmla_f32_z(pg, sum0, a0, b0); sum1 = svmla_f32_z(pg, sum1, a1, b1); - i += 2 * vec_len_f32_sve; + i += 2 * vec_len; } - if (i + vec_len_f32_sve <= n) { + if (i + vec_len <= n) { svfloat32_t a0 = svld1_f32(pg, _a + i); svfloat32_t b0 = svld1_f32(pg, _b + i); sum0 = svmla_f32_z(pg, sum0, a0, b0); - i += vec_len_f32_sve; + i += vec_len; } // Process the tail elements parallely as well (if any) if (i != n) { @@ -701,12 +739,68 @@ static ov::float16 dot_product_fp16(ov::float16* a, float* head_sum) { size_t i = 0; ov::float16 sum = 0.0f; + auto _a = reinterpret_cast(a); + auto _b = reinterpret_cast(b); + +# if defined(HAVE_SVE) + svbool_t pg = svptrue_b16(); + svfloat16_t sum0 = svdup_n_f16(0.0f); + svfloat16_t sum1 = svdup_n_f16(0.0f); + svfloat16_t sum2 = svdup_n_f16(0.0f); + svfloat16_t sum3 = svdup_n_f16(0.0f); + auto vec_len = vec_len_f16_sve(); + + for (; i + 4 * vec_len <= n; i += 4 * vec_len) { + svfloat16_t a0 = svld1_f16(pg, _a + i); + svfloat16_t a1 = svld1_f16(pg, _a + i + vec_len); + svfloat16_t a2 = svld1_f16(pg, _a + i + vec_len * 2); + svfloat16_t a3 = svld1_f16(pg, _a + i + vec_len * 3); + + svfloat16_t b0 = svld1_f16(pg, _b + i); + svfloat16_t b1 = svld1_f16(pg, _b + i + vec_len); + svfloat16_t b2 = svld1_f16(pg, _b + i + vec_len * 2); + svfloat16_t b3 = svld1_f16(pg, _b + i + vec_len * 3); + + sum0 = svmla_f16_z(pg, sum0, a0, b0); + sum1 = svmla_f16_z(pg, sum1, a1, b1); + sum2 = svmla_f16_z(pg, sum2, a2, b2); + sum3 = svmla_f16_z(pg, sum3, a3, b3); + } + if (i + 2 * vec_len <= n) { + svfloat16_t a0 = svld1_f16(pg, _a + i); + svfloat16_t a1 = svld1_f16(pg, _a + i + vec_len); + + svfloat16_t b0 = svld1_f16(pg, _b + i); + svfloat16_t b1 = svld1_f16(pg, _b + i + vec_len); + + sum0 = svmla_f16_z(pg, sum0, a0, b0); + sum1 = svmla_f16_z(pg, sum1, a1, b1); + i += 2 * vec_len; + } + if (i + vec_len <= n) { + svfloat16_t a0 = svld1_f16(pg, _a + i); + svfloat16_t b0 = svld1_f16(pg, _b + i); + sum0 = svmla_f16_z(pg, sum0, a0, b0); + i += vec_len; + } + // Process the tail elements parallely as well (if any) + if (i != n) { + svbool_t pg_rem = svwhilelt_b16(0, static_cast(n - i)); + svfloat16_t a0 = svld1_f16(pg_rem, _a + i); + svfloat16_t b0 = svld1_f16(pg_rem, _b + i); + sum0 = svmla_f16_m(pg_rem, sum0, a0, b0); + i = n; + } + float16_t sum_0 = svaddv_f16(pg, sum0); + float16_t sum_1 = svaddv_f16(pg, sum1); + float16_t sum_2 = svaddv_f16(pg, sum2); + float16_t sum_3 = svaddv_f16(pg, sum3); + sum = static_cast(sum_0 + sum_1 + sum_2 + sum_3); +# else auto vsum0 = vdupq_n_f16(0.0f); auto vsum1 = vdupq_n_f16(0.0f); auto vsum2 = vdupq_n_f16(0.0f); auto vsum3 = vdupq_n_f16(0.0f); - auto _a = reinterpret_cast(a); - auto _b = reinterpret_cast(b); for (; i + 4 * vec_len_f16_neon <= n; i += vec_len_f16_neon * 4) { auto va0 = vld1q_f16(_a + i); @@ -747,7 +841,7 @@ static ov::float16 dot_product_fp16(ov::float16* a, vsum0 = vaddq_f16(vsum0, vsum2); sum = hsum(vsum0); - +# endif for (; i < n; i++) { sum += a[i] * b[i]; } @@ -938,11 +1032,11 @@ static void attn_reduce(T* dst, float* temp, size_t M, size_t S, size_t temp_str #elif defined(OPENVINO_ARCH_ARM64) # if defined(HAVE_SVE) auto _dst = reinterpret_cast(dst); - size_t inc = vec_len_f32_sve; + size_t inc = vec_len_f32_sve(); svbool_t pg = svptrue_b32(); while (i < S) { - if (S - i < vec_len_f32_sve) { + if (S - i < vec_len_f32_sve()) { inc = S - i; pg = svwhilelt_b32(0, static_cast(inc)); } @@ -985,6 +1079,21 @@ static void attn_reduce(T* dst, float* temp, size_t M, size_t S, size_t temp_str #if defined(__ARM_FEATURE_FP16_VECTOR_ARITHMETIC) static void attn_reduce(ov::float16* dst, ov::float16* temp, size_t M, size_t S, size_t temp_stride) { size_t i = 0; +# if defined(HAVE_SVE) + svbool_t pg = svptrue_b16(); + + for (; i + vec_len_f16_sve() <= S; i += vec_len_f16_sve()) { + auto* src = temp + i; + auto result_vec_fp16 = svdup_n_f16(0.0f); + + for (size_t m = 0; m < M; m++) { + auto o_vec_fp16 = svld1_f16(pg, reinterpret_cast(src)); + result_vec_fp16 = svadd_f16_m(pg, result_vec_fp16, o_vec_fp16); + src += temp_stride; + } + svst1_f16(pg, reinterpret_cast(dst + i), result_vec_fp16); + } +# else for (; i + vec_len_f16_neon <= S; i += vec_len_f16_neon) { auto* src = temp + i; auto result_vec_fp16 = vdupq_n_f16(0.0f); @@ -995,6 +1104,7 @@ static void attn_reduce(ov::float16* dst, ov::float16* temp, size_t M, size_t S, } vst1q_f16(reinterpret_cast(dst + i), result_vec_fp16); } +# endif for (; i < S; i++) { auto* src = temp + i; float sum = 0.0f; diff --git a/src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/softmax_kernel.hpp b/src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/softmax_kernel.hpp index 35aab5b59c7d0e..d8eaa7889b9502 100644 --- a/src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/softmax_kernel.hpp +++ b/src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/softmax_kernel.hpp @@ -6,6 +6,7 @@ #include #include #include +#include #include #include "common.hpp" @@ -464,10 +465,61 @@ inline void scale_add2_reduce_max(ov::float16* a, size_t size, float alibi_slope, ov::float16& max) { + size_t i = 0; +# if defined(HAVE_SVE) + svfloat16_t v_max = svdup_n_f16(static_cast(-FLT_MAX)); + svfloat16_t v_scale = svdup_n_f16(static_cast(scale)); + svfloat16_t v_a; + svuint16_t v_zeroi16 = svdup_n_u16(0); + svfloat16_t v_nfltmax = svdup_n_f16(static_cast(-FLT_MAX)); + svfloat16_t v_alibi_slope = svdup_n_f16(static_cast(alibi_slope)); + + svbool_t mask_xor = svptrue_b16(); + if (!select_nfltmax_at_0) + mask_xor = svnot_z(svptrue_b16(), mask_xor); + + svbool_t pg_f16 = svptrue_b16(); + svbool_t pg_u8 = svptrue_b8(); + svbool_t pg_u16 = svptrue_b16(); + size_t inc = vec_len_f16_sve(); + + while (i < size) { + if (size - i < vec_len_f16_sve()) { + inc = size - i; + pg_f16 = svwhilelt_b16(0, static_cast(inc)); + pg_u8 = svwhilelt_b8(0, static_cast(inc)); + pg_u16 = svwhilelt_b16(0, static_cast(inc)); + } + v_a = svld1_f16(pg_f16, reinterpret_cast(a + i)); + v_a = svmul_f16_z(pg_f16, v_a, v_scale); + + if (has_alibi) { + svfloat16_t v_lookup = svld1_f16(pg_f16, reinterpret_cast(alibi_lookup + i)); + v_a = svmla_f16_z(pg_f16, v_a, v_lookup, v_alibi_slope); + } + + if (has_attn_mask) { + svfloat16_t v_mask = svld1_f16(pg_f16, reinterpret_cast(attn_mask + i)); + v_a = svadd_f16_z(pg_f16, v_a, v_mask); + } + + if (has_causal_mask) { + svuint8_t v_maski8 = svld1_u8(pg_u8, causal_mask + i); + svuint16_t v_maski16 = svtrn1_u16(svreinterpret_u16_u8(v_maski8), svdup_n_u16(0)); + svbool_t kmask = svcmpeq_u16(pg_u16, v_maski16, v_zeroi16); + kmask = sveor_z(pg_u16, kmask, mask_xor); + v_a = svsel_f16(kmask, v_nfltmax, v_a); + } + + v_max = svmax_f16_z(pg_f16, v_max, v_a); + svst1_f16(pg_f16, reinterpret_cast(a + i), v_a); + i += inc; + } + max = svmaxv_f16(pg_f16, v_max); +# else float16x8_t v_max = vdupq_n_f16(static_cast(-FLT_MAX)); float16x8_t v_scale = vdupq_n_f16(static_cast(scale)); float16x8_t v_a; - size_t i = 0; uint16x8_t v_zeroi16 = vdupq_n_u16(0); float16x8_t v_nfltmax = vdupq_n_f16(static_cast(-FLT_MAX)); uint16x8_t mask_xor = vdupq_n_u16(select_nfltmax_at_0 ? 0xFFFF : 0); @@ -500,6 +552,7 @@ inline void scale_add2_reduce_max(ov::float16* a, vst1q_f16(reinterpret_cast(a + i), v_a); } max = vmaxvq_f16(v_max); +# endif // process tails for (; i < size; i++) { a[i] *= scale; @@ -664,12 +717,11 @@ inline void exp_reduce_sum(float* a, const float max, const size_t size, float& svfloat32_t v_a; svfloat32_t v_max = svdup_n_f32(max); svfloat32_t v_sum = svdup_n_f32(0.0f); - size_t vec_len_f32_sve = svcntw(); - size_t inc = vec_len_f32_sve; + size_t inc = vec_len_f32_sve(); svbool_t pg = svptrue_b32(); while (i < size) { - if (size - i < vec_len_f32_sve) { + if (size - i < vec_len_f32_sve()) { inc = size - i; pg = svwhilelt_b32(0, static_cast(inc)); } @@ -705,10 +757,47 @@ inline void exp_reduce_sum(float* a, const float max, const size_t size, float& #if defined(OPENVINO_ARCH_ARM64) inline void exp_reduce_sum_f32(ov::float16* a, const ov::float16 max, const size_t size, ov::float16& sum) { + size_t i = 0; +# if defined(HAVE_SVE) + svfloat32_t v_a; + svfloat32_t v_max = svdup_n_f32(static_cast(max)); + svfloat32_t v_sum = svdup_n_f32(0.0f); + + svbool_t pg_f32 = svptrue_b32(); + svbool_t pg_f16 = svptrue_b16(); + svfloat16_t zero = svdup_n_f16(0.0); + size_t inc = vec_len_f32_sve(); + + while (i < size) { + if (size - i < vec_len_f16_sve()) + pg_f16 = svwhilelt_b16(0, static_cast(size - i)); + if (size - i < vec_len_f32_sve()) { + pg_f32 = svwhilelt_b32(0, static_cast(size - i)); + inc = size - i; + } + // Load 16 elements and interleave with zeros so we have 8 elements with 0 in high parts + svfloat16_t v_a_f16 = svld1_f16(pg_f16, reinterpret_cast(a + i)); + v_a_f16 = svzip1_f16(v_a_f16, zero); + + // Convert to f32 and perform required operations + v_a = svcvt_f32_f16_z(pg_f16, v_a_f16); + v_a = svsub_f32_z(pg_f32, v_a, v_max); + v_a = exp_ps_sve(pg_f32, v_a); + v_sum = svadd_f32_z(pg_f32, v_sum, v_a); + + // Convert to f16 and compact non-zero elements (even indices) to the low part + // so that we can store them in the result using svwhilelt + svfloat16_t v_result = svcvt_f16_f32_z(pg_f32, v_a); + v_result = svtbl_f16(v_result, svindex_u16(0, 2)); + + svst1_f16(svwhilelt_b16(0, static_cast(inc)), reinterpret_cast(a + i), v_result); + i += inc; + } + float total_sum = svaddv_f32(svptrue_b32(), v_sum); +# else float32x4_t v_a; float32x4_t v_max = vdupq_n_f32(static_cast(max)); float32x4_t v_sum = vdupq_n_f32(0.0f); - size_t i = 0; // Process 4 FP32 elements at a time for (; i + vec_len_f32_neon <= size; i += vec_len_f32_neon) { @@ -728,7 +817,7 @@ inline void exp_reduce_sum_f32(ov::float16* a, const ov::float16 max, const size // Reduce sum float total_sum = vaddvq_f32(v_sum); - +# endif // Handle remaining elements for (; i < size; ++i) { float val = exp(static_cast(a[i] - max)); @@ -742,11 +831,32 @@ inline void exp_reduce_sum_f32(ov::float16* a, const ov::float16 max, const size #if defined(__ARM_FEATURE_FP16_VECTOR_ARITHMETIC) inline void exp_reduce_sum(ov::float16* a, const ov::float16 max, const size_t size, ov::float16& sum) { + size_t i = 0; +# if defined(HAVE_SVE) + svfloat16_t v_a; + svfloat16_t v_max = svdup_n_f16(max); + svfloat16_t v_sum = svdup_n_f16(0.0f); + svbool_t pg = svptrue_b16(); + size_t inc = vec_len_f16_sve(); + + while (i < size) { + if (size - i < vec_len_f16_sve()) { + inc = size - i; + pg = svwhilelt_b16(0, static_cast(inc)); + } + v_a = svld1_f16(pg, reinterpret_cast(a + i)); + v_a = svsub_f16_z(pg, v_a, v_max); + v_a = exp_ps_sve_f16(pg, v_a); + v_sum = svadd_f16_m(pg, v_sum, v_a); + svst1_f16(pg, reinterpret_cast(a + i), v_a); + i += inc; + } + sum = svaddv_f16(svptrue_b16(), v_sum); +# else const size_t vec_len_f16_neon = 8; float16x8_t v_a; float16x8_t v_max = vdupq_n_f16(max); float16x8_t v_sum = vdupq_n_f16(0.0f); - size_t i = 0; for (; i + vec_len_f16_neon <= size; i += vec_len_f16_neon) { v_a = vld1q_f16(reinterpret_cast(a + i)); @@ -761,7 +871,7 @@ inline void exp_reduce_sum(ov::float16* a, const ov::float16 max, const size_t s float16x4_t v_sum_final = vpadd_f16(v_sum_pair, v_sum_pair); sum += vget_lane_f16(v_sum_final, 0); - +# endif for (; i < size; ++i) { a[i] = static_cast(exp(static_cast(a[i] - max))); sum += a[i]; @@ -808,11 +918,11 @@ inline void multiply_scalar(float* a, float* a_dst, const float val, const size_ #elif defined(OPENVINO_ARCH_ARM64) # if defined(HAVE_SVE) svfloat32_t v_scale = svdup_n_f32(val); - size_t inc = vec_len_f32_sve; + size_t inc = vec_len_f32_sve(); svbool_t pg = svptrue_b32(); while (i < size) { - if (size - i < vec_len_f32_sve) { + if (size - i < vec_len_f32_sve()) { inc = size - i; pg = svwhilelt_b32(0, static_cast(inc)); } @@ -864,6 +974,7 @@ inline void multiply_scalar(float* a, T* a_dst, const float val, const size_t si } #endif } + #if defined(OPENVINO_ARCH_ARM64) inline void multiply_scalar(ov::float16* a, float* a_dst, const ov::float16 val, const size_t size) { float16x4_t v_a_f16; @@ -902,21 +1013,40 @@ inline void multiply_scalar_f32(ov::float16* a, ov::float16* a_dst, const ov::fl } } #endif + #if defined(__ARM_FEATURE_FP16_VECTOR_ARITHMETIC) inline void multiply_scalar(ov::float16* a, ov::float16* a_dst, const ov::float16 val, const size_t size) { + size_t i = 0; +# if defined(HAVE_SVE) + svfloat16_t v_scale = svdup_n_f16(val); + size_t inc = vec_len_f16_sve(); + svbool_t pg = svptrue_b16(); + + while (i < size) { + if (size - i < vec_len_f16_sve()) { + inc = size - i; + pg = svwhilelt_b16(0, static_cast(inc)); + } + svfloat16_t v_a = svld1_f16(pg, reinterpret_cast(a + i)); + v_a = svmul_f16_z(pg, v_a, v_scale); + svst1_f16(pg, reinterpret_cast(a_dst + i), v_a); + i += inc; + } +# else float16x8_t v_a, v_res; float16x8_t v_val = vdupq_n_f16(val); - size_t i = 0; for (; i + vec_len_f16_neon <= size; i += vec_len_f16_neon) { v_a = vld1q_f16(reinterpret_cast(a + i)); v_res = vmulq_f16(v_a, v_val); vst1q_f16(reinterpret_cast(a_dst + i), v_res); } +# endif for (; i < size; ++i) { a_dst[i] = a[i] * val; } } #endif + template inline void attn_softmax_kernel(T* a, void* a_dst, diff --git a/src/plugins/intel_cpu/src/nodes/kernels/x64/brgemm_kernel.cpp b/src/plugins/intel_cpu/src/nodes/kernels/x64/brgemm_kernel.cpp index 7df2e2371a843a..5ef9d8df003a86 100644 --- a/src/plugins/intel_cpu/src/nodes/kernels/x64/brgemm_kernel.cpp +++ b/src/plugins/intel_cpu/src/nodes/kernels/x64/brgemm_kernel.cpp @@ -289,7 +289,7 @@ void BrgemmKernel::init_brgemm_copy_b( brgemm_matmul_conf_t brgCopyKernelConf; brgCopyKernelConf.src_dt = is_avx_f16_only ? dnnl_data_type_t::dnnl_f32 : dt_in0; brgCopyKernelConf.wei_dt = is_avx_f16_only ? dnnl_data_type_t::dnnl_f32 : dt_in1; - brgCopyKernelConf.orig_wei_dt = dt_in1; + brgCopyKernelConf.orig_wei_dt = static_cast(DnnlExtensionUtils::ElementTypeToDataType(inType)); brgCopyKernelConf.wei_n_blk = N_blk; brgCopyKernelConf.wei_tag = transpose ? dnnl_ba : dnnl_ab; brgCopyKernelConf.copy_B_wei_stride = copy_B_wei_stride; diff --git a/src/plugins/intel_cpu/src/nodes/kernels/x64/mlp_kernel.cpp b/src/plugins/intel_cpu/src/nodes/kernels/x64/mlp_kernel.cpp index 089461bc504172..16f41f7476249e 100644 --- a/src/plugins/intel_cpu/src/nodes/kernels/x64/mlp_kernel.cpp +++ b/src/plugins/intel_cpu/src/nodes/kernels/x64/mlp_kernel.cpp @@ -550,12 +550,13 @@ void GateUpCombine::generate() { const auto zmm_up = zmm0; const auto ymm_dst = ymm5; - auto injector = std::make_shared>( + auto injector = std::make_shared>( this, m_act_alg, 1.f, 1.0f, 1.f, + data_type::f32, true, // save_state, true due to additional r15 is used. Xbyak::Reg64(Xbyak::Operand::R10), // p_table Xbyak::Opmask(1), // k_mask diff --git a/src/plugins/intel_cpu/src/nodes/kernels/x64/non_max_suppression.cpp b/src/plugins/intel_cpu/src/nodes/kernels/x64/non_max_suppression.cpp index 92c1c895ffd19e..ab5d89f3d68c48 100644 --- a/src/plugins/intel_cpu/src/nodes/kernels/x64/non_max_suppression.cpp +++ b/src/plugins/intel_cpu/src/nodes/kernels/x64/non_max_suppression.cpp @@ -6,6 +6,7 @@ #include "utils/general_utils.h" +using namespace dnnl::impl; using namespace dnnl::impl::cpu; #define GET_OFF(field) offsetof(NmsCallArgs, field) @@ -20,7 +21,7 @@ void NonMaxSuppression::generate() { load_scalar_emitter.reset(new jit_load_emitter(this, isa, ov::element::f32, ov::element::f32, scalar_step)); exp_injector.reset( - new x64::jit_uni_eltwise_injector_f32(this, dnnl::impl::alg_kind::eltwise_exp, 0.f, 0.f, 1.f)); + new x64::jit_uni_eltwise_injector(this, dnnl::impl::alg_kind::eltwise_exp, 0.f, 0.f, 1.f, data_type::f32)); this->preamble(); diff --git a/src/plugins/intel_cpu/src/nodes/kernels/x64/non_max_suppression.hpp b/src/plugins/intel_cpu/src/nodes/kernels/x64/non_max_suppression.hpp index a47120de5a9040..a489faf9517bfa 100644 --- a/src/plugins/intel_cpu/src/nodes/kernels/x64/non_max_suppression.hpp +++ b/src/plugins/intel_cpu/src/nodes/kernels/x64/non_max_suppression.hpp @@ -110,7 +110,7 @@ class NonMaxSuppression : public JitKernel { Xbyak::Opmask k_mask = Xbyak::Opmask(7); Xbyak::Opmask k_mask_one = Xbyak::Opmask(6); - std::shared_ptr> exp_injector; + std::shared_ptr> exp_injector; inline void hard_nms(); diff --git a/src/plugins/intel_cpu/src/nodes/mvn.cpp b/src/plugins/intel_cpu/src/nodes/mvn.cpp index 6c11feee2534b9..bab6a3ed1b8091 100644 --- a/src/plugins/intel_cpu/src/nodes/mvn.cpp +++ b/src/plugins/intel_cpu/src/nodes/mvn.cpp @@ -949,11 +949,12 @@ struct jit_uni_mvn_kernel_f32 : public jit_uni_mvn_kernel, public jit_generator for (int i = 0; i < p.len(); i++) { auto& post_op = p.entry_[i]; if (post_op.is_eltwise()) { - eltwise_injectors.push_back(std::make_shared>(this, - post_op.eltwise.alg, - post_op.eltwise.alpha, - post_op.eltwise.beta, - post_op.eltwise.scale)); + eltwise_injectors.push_back(std::make_shared>(this, + post_op.eltwise.alg, + post_op.eltwise.alpha, + post_op.eltwise.beta, + post_op.eltwise.scale, + data_type::f32)); } else if (post_op.is_depthwise()) { depthwise_injectors.push_back(std::make_shared>(this, post_op)); } else if (post_op.is_quantization()) { @@ -1093,7 +1094,7 @@ struct jit_uni_mvn_kernel_f32 : public jit_uni_mvn_kernel, public jit_generator const int tile_size[kTileNum] = {8, 4, 2, 1}; - std::vector>> eltwise_injectors; + std::vector>> eltwise_injectors; std::vector>> depthwise_injectors; std::vector>> quantization_injectors; diff --git a/src/plugins/intel_cpu/src/nodes/normalize.cpp b/src/plugins/intel_cpu/src/nodes/normalize.cpp index 45e4d12c7e76c6..e6a730027929a9 100644 --- a/src/plugins/intel_cpu/src/nodes/normalize.cpp +++ b/src/plugins/intel_cpu/src/nodes/normalize.cpp @@ -224,11 +224,12 @@ struct jit_uni_normalize_kernel_f32 : public jit_uni_normalize_kernel, public ji for (int i = 0; i < p.len(); i++) { auto& post_op = p.entry_[i]; if (post_op.is_eltwise()) { - eltwise_injectors.push_back(std::make_shared>(this, - post_op.eltwise.alg, - post_op.eltwise.alpha, - post_op.eltwise.beta, - post_op.eltwise.scale)); + eltwise_injectors.push_back(std::make_shared>(this, + post_op.eltwise.alg, + post_op.eltwise.alpha, + post_op.eltwise.beta, + post_op.eltwise.scale, + data_type::f32)); } else if (post_op.is_depthwise()) { depthwise_injectors.push_back(std::make_shared>(this, post_op)); } else if (post_op.is_quantization()) { @@ -310,7 +311,7 @@ struct jit_uni_normalize_kernel_f32 : public jit_uni_normalize_kernel, public ji std::unique_ptr uni_vcvtneps2bf16 = nullptr; - std::vector>> eltwise_injectors; + std::vector>> eltwise_injectors; std::vector>> depthwise_injectors; std::vector>> quantization_injectors; diff --git a/src/plugins/intel_cpu/src/nodes/reduce.cpp b/src/plugins/intel_cpu/src/nodes/reduce.cpp index 7828e11f9d413a..3a209ac132c560 100644 --- a/src/plugins/intel_cpu/src/nodes/reduce.cpp +++ b/src/plugins/intel_cpu/src/nodes/reduce.cpp @@ -142,8 +142,12 @@ struct jit_uni_reduce_kernel_f32 : public jit_uni_reduce_kernel, public jit_gene void generate() override { if (jcp_.reduce_mode == Algorithm::ReduceLogSumExp) { - exp_injector = - std::make_shared>(this, alg_kind::eltwise_exp, 0.f, 0.f, 1.f); + exp_injector = std::make_shared>(this, + alg_kind::eltwise_exp, + 0.f, + 0.f, + 1.f, + data_type::f32); } if (mayiuse(avx512_core)) @@ -244,7 +248,7 @@ struct jit_uni_reduce_kernel_f32 : public jit_uni_reduce_kernel, public jit_gene Xbyak::Label l_table; std::shared_ptr uni_vcvtneps2bf16; - std::shared_ptr> exp_injector; + std::shared_ptr> exp_injector; inline void reduce_main() { // ================================================================ @@ -1206,11 +1210,12 @@ struct jit_uni_reduce_post_kernel_f32 : public jit_uni_reduce_post_kernel, publi for (int i = 0; i < p.len(); i++) { auto& post_op = p.entry_[i]; if (post_op.is_eltwise()) { - eltwise_injectors.push_back(std::make_shared>(this, - post_op.eltwise.alg, - post_op.eltwise.alpha, - post_op.eltwise.beta, - post_op.eltwise.scale)); + eltwise_injectors.push_back(std::make_shared>(this, + post_op.eltwise.alg, + post_op.eltwise.alpha, + post_op.eltwise.beta, + post_op.eltwise.scale, + data_type::f32)); } else if (post_op.is_depthwise()) { depthwise_injectors.push_back(std::make_shared>(this, post_op)); } else if (post_op.is_quantization()) { @@ -1224,8 +1229,12 @@ struct jit_uni_reduce_post_kernel_f32 : public jit_uni_reduce_post_kernel, publi } if (jcp_.reduce_mode == Algorithm::ReduceLogSum || jcp_.reduce_mode == Algorithm::ReduceLogSumExp) { - log_injector = - std::make_shared>(this, alg_kind::eltwise_log, 0.f, 0.f, 1.f); + log_injector = std::make_shared>(this, + alg_kind::eltwise_log, + 0.f, + 0.f, + 1.f, + data_type::f32); } if (mayiuse(avx512_core)) @@ -1336,9 +1345,9 @@ struct jit_uni_reduce_post_kernel_f32 : public jit_uni_reduce_post_kernel, publi Vmm vmm_d_bias = Vmm(8); std::shared_ptr uni_vcvtneps2bf16; - std::shared_ptr> log_injector; + std::shared_ptr> log_injector; - std::vector>> eltwise_injectors; + std::vector>> eltwise_injectors; std::vector>> depthwise_injectors; std::vector>> quantization_injectors; diff --git a/src/plugins/intel_cpu/src/nodes/region_yolo.cpp b/src/plugins/intel_cpu/src/nodes/region_yolo.cpp index b7d98064713fd0..56f600c7f8ac29 100644 --- a/src/plugins/intel_cpu/src/nodes/region_yolo.cpp +++ b/src/plugins/intel_cpu/src/nodes/region_yolo.cpp @@ -18,6 +18,7 @@ #include "openvino/opsets/opset1.hpp" #include "utils/bfloat16.hpp" +using namespace dnnl::impl; using namespace dnnl::impl::cpu; using namespace dnnl::impl::cpu::x64; using namespace dnnl::impl::utils; @@ -46,7 +47,7 @@ struct jit_uni_logistic_kernel_f32 : public jit_uni_logistic_kernel, public jit_ void generate() override { exp_injector.reset( - new jit_uni_eltwise_injector_f32(this, dnnl::impl::alg_kind::eltwise_exp, 0.f, 0.f, 1.f)); + new jit_uni_eltwise_injector(this, dnnl::impl::alg_kind::eltwise_exp, 0.f, 0.f, 1.f, data_type::f32)); if (mayiuse(avx512_core)) uni_vcvtneps2bf16.reset(new jit_uni_vcvtneps2bf16(this, isa)); @@ -134,7 +135,7 @@ struct jit_uni_logistic_kernel_f32 : public jit_uni_logistic_kernel, public jit_ Xbyak::Label l_table; - std::shared_ptr> exp_injector; + std::shared_ptr> exp_injector; jit_logistic_config_params jcp_; diff --git a/src/plugins/intel_cpu/src/nodes/subgraph.cpp b/src/plugins/intel_cpu/src/nodes/subgraph.cpp index 2b0c7b55fb043d..ba6f1eda215dce 100644 --- a/src/plugins/intel_cpu/src/nodes/subgraph.cpp +++ b/src/plugins/intel_cpu/src/nodes/subgraph.cpp @@ -236,7 +236,7 @@ struct SubgraphCodeGeneratorKey { bool operator==(const SubgraphCodeGeneratorKey& rhs) const; std::shared_ptr attrs = nullptr; - uint8_t broadcasting_mask = 0; + uint32_t broadcasting_mask = 0; }; struct SubgraphShapeInferResultKey { @@ -738,8 +738,10 @@ Subgraph::ControlFlowPasses Subgraph::getControlFlowPasses() const { return backend_passes; } -uint8_t Subgraph::getBroadcastingMask(const std::vector& input_shapes) { - uint8_t mask = 0; +uint32_t Subgraph::getBroadcastingMask(const std::vector& input_shapes) { + uint32_t mask = 0; + OPENVINO_ASSERT(broadcastable_inputs.size() <= sizeof(mask) * CHAR_BIT, + "Incorrect size of broadcastable inputs of Subgraph"); for (const auto& broadcastable_input : broadcastable_inputs) { const auto& shape = input_shapes[broadcastable_input.first]; mask = mask << 1; diff --git a/src/plugins/intel_cpu/src/nodes/subgraph.h b/src/plugins/intel_cpu/src/nodes/subgraph.h index aac0fa1ea2f535..a84e46d9ae02da 100644 --- a/src/plugins/intel_cpu/src/nodes/subgraph.h +++ b/src/plugins/intel_cpu/src/nodes/subgraph.h @@ -70,7 +70,7 @@ class Subgraph : public Node { std::pair, std::vector> getIOPrecisions() const; static uint64_t getBodyHash(const std::shared_ptr& snippet); - uint8_t getBroadcastingMask(const std::vector& input_shapes); + uint32_t getBroadcastingMask(const std::vector& input_shapes); using DataFlowPasses = std::vector; using ControlFlowPasses = std::vector; diff --git a/src/plugins/intel_cpu/src/transformations/snippets/x64/pass/lowered/insert_brgemm_copy_buffers.cpp b/src/plugins/intel_cpu/src/transformations/snippets/x64/pass/lowered/insert_brgemm_copy_buffers.cpp index b2bab35a17c1ad..d68469e0bde486 100644 --- a/src/plugins/intel_cpu/src/transformations/snippets/x64/pass/lowered/insert_brgemm_copy_buffers.cpp +++ b/src/plugins/intel_cpu/src/transformations/snippets/x64/pass/lowered/insert_brgemm_copy_buffers.cpp @@ -58,6 +58,7 @@ bool InsertBrgemmCopyBuffers::run(LinearIR& linear_ir, LinearIR::constExprIt beg const auto vnni_factor = brgemm_utils::compute_vnni_factor(src_dt); const auto inner_k_blk = brgemm_utils::repacking::compute_inner_k_block(src_dt); + OPENVINO_ASSERT(inner_k_blk > 0, "inner_k_blk cannot be zero!"); const auto tile_scratch_size = BrgemmCPU::SCRATCH_BYTE_SIZE; const auto current_scratch_size = scratch_expr->get_byte_size(); OPENVINO_ASSERT(current_scratch_size == tile_scratch_size, diff --git a/src/plugins/intel_cpu/src/transformations/transformation_pipeline.cpp b/src/plugins/intel_cpu/src/transformations/transformation_pipeline.cpp index 4d7df9a335e98a..81eb70d328630d 100644 --- a/src/plugins/intel_cpu/src/transformations/transformation_pipeline.cpp +++ b/src/plugins/intel_cpu/src/transformations/transformation_pipeline.cpp @@ -12,6 +12,9 @@ #include #include "openvino/op/paged_attention.hpp" +#include "openvino/op/prelu.hpp" +#include "openvino/op/round.hpp" +#include "openvino/op/sqrt.hpp" #include "openvino/opsets/opset1.hpp" #include "openvino/opsets/opset10.hpp" #include "openvino/opsets/opset2.hpp" @@ -1123,14 +1126,17 @@ void Transformations::MainSnippets(void) { ov::is_type(n) || ov::is_type(n) || ov::is_type(n) || ov::is_type(n) || ov::is_type(n) || ov::is_type(n) || - ov::is_type(n) || ov::is_type(n) || - ov::is_type(n) || ov::is_type(n) || - ov::is_type(n) || ov::is_type(n) || + ov::is_type(n) || ov::is_type(n) || + ov::is_type(n) || ov::is_type(n) || + ov::is_type(n) || ov::is_type(n) || + ov::is_type(n) || ov::is_type(n) || + ov::is_type(n) || ov::is_type(n) || ov::is_type(n) || ov::is_type(n) || ov::is_type(n) || ov::is_type(n) || - ov::is_type(n) || ov::is_type(n) || - ov::is_type(n) || ov::is_type(n) || - ov::is_type(n)); + ov::is_type(n) || ov::is_type(n) || + ov::is_type(n) || ov::is_type(n) || + ov::is_type(n) || ov::is_type(n) || + ov::is_type(n) || ov::is_type(n)); #else // CPU Plugin support Swish in Subgraph via conversion to SwichCPU which assumes second input to be constant, // and CPU Plugin does not support Mish for x64 diff --git a/src/plugins/intel_cpu/src/utils/verbose.cpp b/src/plugins/intel_cpu/src/utils/verbose.cpp index 3e5bb7e8ef7930..c9283a130e3d6e 100644 --- a/src/plugins/intel_cpu/src/utils/verbose.cpp +++ b/src/plugins/intel_cpu/src/utils/verbose.cpp @@ -107,7 +107,7 @@ void Verbose::printInfo() { shift(written); written = snprintf(portsInfo + written_total, CPU_VERBOSE_DAT_LEN - written_total, "%s", prefix.c_str()); shift(written); - std::string fmt_str = dnnl::impl::md2fmt_str(desc, dnnl::impl::format_kind_t::dnnl_format_kind_undef); + std::string fmt_str = dnnl::impl::md2fmt_str("", desc, dnnl::impl::format_kind_t::dnnl_format_kind_undef); written = snprintf(portsInfo + written_total, CPU_VERBOSE_DAT_LEN - written_total, "%s", fmt_str.c_str()); shift(written); written = snprintf(portsInfo + written_total, CPU_VERBOSE_DAT_LEN - written_total, ":"); diff --git a/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/classes/activation.cpp b/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/classes/activation.cpp index 0f63a7517b5745..16ce7ff630c20e 100644 --- a/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/classes/activation.cpp +++ b/src/plugins/intel_cpu/tests/functional/custom/single_layer_tests/classes/activation.cpp @@ -197,7 +197,9 @@ std::string ActivationLayerCPUTest::getPrimitiveType(const utils::ActivationType (activation_type == utils::ActivationTypes::Sqrt) || (activation_type == utils::ActivationTypes::Swish) || (activation_type == utils::ActivationTypes::LogicalNot) || - (activation_type == utils::ActivationTypes::Tanh))) { + (activation_type == utils::ActivationTypes::Tanh) || + (activation_type == utils::ActivationTypes::RoundHalfAwayFromZero) || + (activation_type == utils::ActivationTypes::RoundHalfToEven))) { return "jit"; } @@ -209,7 +211,9 @@ std::string ActivationLayerCPUTest::getPrimitiveType(const utils::ActivationType if ((activation_type == utils::ActivationTypes::Floor) || (activation_type == utils::ActivationTypes::Ceiling) || (activation_type == utils::ActivationTypes::IsNaN) || - (activation_type == utils::ActivationTypes::IsFinite)) { + (activation_type == utils::ActivationTypes::IsFinite) || + (activation_type == utils::ActivationTypes::RoundHalfAwayFromZero) || + (activation_type == utils::ActivationTypes::RoundHalfToEven)) { return "ref"; } return "acl"; @@ -265,22 +269,26 @@ const std::map>>& activat const std::map>>& activationTypesSnippets() { static const std::map>> activationTypes { - {Abs, {{}}}, - {Exp, {{}}}, - {Ceiling, {{}}}, - {Clamp, {{-2.0f, 2.0f}}}, - {Elu, {{0.1f}}}, - {Floor, {{}}}, - {GeluErf, {{}}}, - {GeluTanh, {{}}}, - {Relu, {{}}}, - {HSwish, {{}}}, + {Abs, {{}}}, + {Exp, {{}}}, + {Ceiling, {{}}}, + {Clamp, {{-2.0f, 2.0f}}}, + {Elu, {{0.1f}}}, + {Floor, {{}}}, + {GeluErf, {{}}}, + {GeluTanh, {{}}}, + {Relu, {{}}}, + {HSwish, {{}}}, + {PReLu, {{-0.01f}}}, + {Sqrt, {{}}}, + {RoundHalfToEven, {{}}}, + {RoundHalfAwayFromZero, {{}}}, #if defined(OPENVINO_ARCH_ARM64) - {Mish, {{}}}, + {Mish, {{}}}, #endif - {Sigmoid, {{}}}, - {Swish, {{0.1f}}}, - {Tanh, {{}}}, + {Sigmoid, {{}}}, + {Swish, {{0.1f}}}, + {Tanh, {{}}}, }; return activationTypes; 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 4c34b3fd2506ac..a88ba297dc57ee 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 @@ -512,12 +512,18 @@ std::vector disabledTestPatterns() { R"(.*smoke_MatMulCompressedWeights_corner_cases_basic/MatmulWeightsDecompression.CompareWithRefs/data_shape=\[\?.\?.\?\]_\(\[1,1,4096\]\)_weights_shape=\[4096,4096\]_group_size=128_weights_precision=nf4_decompression_precision=f16_scale_precision=undefined_transpose_weights=0_decompression_subtract=full_reshape_on_decompression=1_config=\(\).*)"); retVector.emplace_back(R"(.*smoke_RDFT_CPU_1D/RDFTTestCPU.CompareWithRefs/prec=f32_IS0=\[\]_TS0=\(\(126\)\)_constAxes=true_axes=\(\(0\)\)_isInverse=false.*)"); retVector.emplace_back(R"(.*smoke_RDFT_CPU_2D/RDFTTestCPU.CompareWithRefs/prec=f32_IS0=\[\]_TS0=\(\(16.38\)\)_constAxes=true_axes=\(\(0.1\)\)_isInverse=false.*)"); + // Issue: MFDNN-12818 + retVector.emplace_back(R"(.*smoke_LPT/RecurrentCellTransformation.CompareWithRefImpl/f32_\[1,1,3\]_CPU_f32FQ_X_level=256_.*_FQ_W_level=255.*)"); + retVector.emplace_back(R"(.*smoke_static/ConvertFqRnnToQuantizedRnn.CompareWithRefs/Type=GRUSequence.*2.5.10.*2.1.4.*2.1.4.*)"); #endif if (!ov::with_cpu_x86_avx512_core()) { // on platforms which do not support bfloat16, we are disabling bf16 tests since there are no bf16 primitives, // tests are useless on such platforms retVector.emplace_back(R"(.*(BF|bf)16.*)"); retVector.emplace_back(R"(.*bfloat16.*)"); + // Issue: MFDNN-12818 + retVector.emplace_back(R"(.*smoke_LPT/RecurrentCellTransformation.CompareWithRefImpl/f32_\[1,1,3\]_CPU_f32FQ_X_level=256_.*_FQ_W_level=255.*)"); + retVector.emplace_back(R"(.*smoke_static/ConvertFqRnnToQuantizedRnn.CompareWithRefs/Type=GRUSequence.*2.5.10.*2.1.4.*2.1.4.*)"); } if (!ov::with_cpu_x86_avx2()) { // MatMul in Snippets uses BRGEMM that is supported only on AVX2 (and newer) platforms diff --git a/src/plugins/intel_cpu/thirdparty/onednn b/src/plugins/intel_cpu/thirdparty/onednn index c60a9946aa2386..ec3d68916fb1b0 160000 --- a/src/plugins/intel_cpu/thirdparty/onednn +++ b/src/plugins/intel_cpu/thirdparty/onednn @@ -1 +1 @@ -Subproject commit c60a9946aa2386890e5c9f5587974facb7624227 +Subproject commit ec3d68916fb1b03554e00b35b568874ba56093b6 diff --git a/src/plugins/intel_gpu/include/intel_gpu/plugin/common_utils.hpp b/src/plugins/intel_gpu/include/intel_gpu/plugin/common_utils.hpp index 792745193ed550..c958b840e65290 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/plugin/common_utils.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/plugin/common_utils.hpp @@ -6,7 +6,6 @@ #include #include -#include "intel_gpu/runtime/engine.hpp" #include "intel_gpu/runtime/layout.hpp" #include "intel_gpu/runtime/memory.hpp" #include "intel_gpu/runtime/optionals.hpp" @@ -104,26 +103,6 @@ inline ov::Shape predict_shape(const std::string& name, const cldnn::layout layo return layout.get_shape(); } -inline cldnn::memory::ptr allocate_memory_evenif_zero_bytes(cldnn::engine& _engine, - const cldnn::layout& layout, - cldnn::allocation_type type, - bool reset = true) { - if (layout.bytes_count() == 0) { - auto non_zero_layout = cldnn::layout({1}, layout.data_type, layout.format); - auto res = _engine.allocate_memory(non_zero_layout, type, false); - return _engine.reinterpret_buffer(*res, layout); - } else { - return _engine.allocate_memory(layout, type, reset); - } -} - -inline cldnn::memory::ptr allocate_memory_evenif_zero_bytes(cldnn::engine& _engine, - const cldnn::layout& layout, - bool reset = true) { - cldnn::allocation_type type = _engine.get_lockable_preferred_memory_allocation_type(layout.format.is_image_2d()); - return allocate_memory_evenif_zero_bytes(_engine, layout, type, reset); -} - /// WA: Force exit. Any opencl api call can be hang after CL_OUT_OF_RESOURCES. inline void ForceExit() { std::cerr << "[GPU] force exit.\n" diff --git a/src/plugins/intel_gpu/include/intel_gpu/runtime/execution_config.hpp b/src/plugins/intel_gpu/include/intel_gpu/runtime/execution_config.hpp index 5e059b17da0e97..70a04f0b0c3a99 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/runtime/execution_config.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/runtime/execution_config.hpp @@ -150,6 +150,7 @@ class ExecutionConfig { void apply_performance_hints(const cldnn::device_info& info); void apply_priority_hints(const cldnn::device_info& info); void apply_debug_options(const cldnn::device_info& info); + void update_specific_default_properties(const cldnn::device_info& info); template void apply_rt_info_property(const ov::Property& property, const ov::RTMap& rt_info) { @@ -167,6 +168,8 @@ class ExecutionConfig { std::map supported_properties; std::map property_validators; + + bool specific_default_properties_is_set = false; }; } // namespace intel_gpu diff --git a/src/plugins/intel_gpu/src/graph/loop.cpp b/src/plugins/intel_gpu/src/graph/loop.cpp index 5d842c2a863433..f381cf3638e022 100644 --- a/src/plugins/intel_gpu/src/graph/loop.cpp +++ b/src/plugins/intel_gpu/src/graph/loop.cpp @@ -7,7 +7,6 @@ #include "mutable_data_inst.h" #include "json_object.h" #include "primitive_type_base.h" -#include "intel_gpu/plugin/common_utils.hpp" #include "intel_gpu/primitives/data.hpp" #include "intel_gpu/primitives/mutable_data.hpp" #include @@ -319,7 +318,7 @@ void loop_inst::update_backedge_mapped_memory() { // generally, shouldn't go this way, but... auto output_prim = body_network->get_primitive(back_edge.from); layout output_layout = output_prim->output_memory().get_layout(); - backedge_mem = ov::intel_gpu::allocate_memory_evenif_zero_bytes(body_network->get_engine(), output_layout, false); + backedge_mem = body_network->get_engine().allocate_memory(output_layout, 0); } } else { auto external_id = output_mapping.front()->external_id; @@ -397,7 +396,7 @@ loop_inst::concatenated_memory_mapping::ptr loop_inst::create_concat_memory_map( << sliced_layout.get_partial_shape().to_string() << " to " << updated_sliced_layout.to_string() << std::endl; sliced_layout.set_partial_shape(updated_sliced_layout); - inter_mem_ptr = ov::intel_gpu::allocate_memory_evenif_zero_bytes(engine, sliced_layout); + inter_mem_ptr = engine.allocate_memory(sliced_layout); intern_prim->set_output_layout(sliced_layout, internal_id.idx); } @@ -407,8 +406,8 @@ loop_inst::concatenated_memory_mapping::ptr loop_inst::create_concat_memory_map( } else { sliced_mems.reserve(num_iterations); sliced_mems.push_back(inter_mem_ptr); - for (int j=1; j < num_iterations; ++j) { - memory::ptr sliced_mem = ov::intel_gpu::allocate_memory_evenif_zero_bytes(engine, sliced_layout); + for (int j = 1; j < num_iterations; ++j) { + memory::ptr sliced_mem = engine.allocate_memory(sliced_layout); sliced_mems.push_back(sliced_mem); } } @@ -499,7 +498,7 @@ void loop_inst::preprocess_input_memory(const int64_t num_iterations) { // if internal input memory is in backedge, allocate new memory. // Because internal input memory's data will be updated through backedge process. if (iter != _back_edges.end()) { - internal_input_memory = ov::intel_gpu::allocate_memory_evenif_zero_bytes(body_network->get_engine(), memory->get_layout(), false); + internal_input_memory = body_network->get_engine().allocate_memory(memory->get_layout(), false); internal_input_memory->copy_from(body_network->get_stream(), *memory); GPU_DEBUG_LOG << "Input memory of internal node(" << internal_id.to_string() << ") is set to new memory(" << internal_input_memory << ", " << internal_input_memory->get_layout().to_short_string() @@ -722,7 +721,7 @@ void loop_inst::postprocess_output_memory(bool is_dynamic, int64_t current_itera } else { if (!output_allocated || get_flag(ExecutionFlags::SHAPE_CHANGED)) { auto concat_layout = _impl_params->get_output_layout(external_id.idx); - auto concat_mem = ov::intel_gpu::allocate_memory_evenif_zero_bytes(_network.get_engine(), concat_layout, false); + auto concat_mem = _network.get_engine().allocate_memory(concat_layout, false); external_outputs[external_id.idx] = concat_mem; auto iter = std::find_if(concatenated_output_mem_mappings.begin(), concatenated_output_mem_mappings.end(), @@ -1081,7 +1080,7 @@ std::vector loop_inst::handle_buffers_for_next_iteration(const loop_ // Check backedge_to shape needs to be updated by initial_mem OPENVINO_ASSERT(mapping.initial_mem != nullptr, "initial_mem should not be null"); if (!mapping.initial_mem->get_layout().identical(to_mem->get_layout())) { - to_mem = ov::intel_gpu::allocate_memory_evenif_zero_bytes(body_network->get_engine(), mapping.initial_mem->get_layout(), false); + to_mem = body_network->get_engine().allocate_memory(mapping.initial_mem->get_layout(), false); body_network->set_input_data(to_id, to_mem); ev = to_mem->copy_from(body_network->get_stream(), *(mapping.initial_mem)); @@ -1104,7 +1103,7 @@ std::vector loop_inst::handle_buffers_for_next_iteration(const loop_ // Check backedge_to shape needs to be updated by backedge_from if (!from_mem->get_layout().identical(to_mem->get_layout())) { - to_mem = ov::intel_gpu::allocate_memory_evenif_zero_bytes(body_network->get_engine(), from_mem->get_layout(), false); + to_mem = body_network->get_engine().allocate_memory(from_mem->get_layout(), false); GPU_DEBUG_LOG << iter << ") [SINGLE] Backedge_to node(" << to_id << ") is set to new memory(" << to_mem << ", " << to_mem->get_layout().to_short_string() << ") because of shape update from backedge_from()" << from_id diff --git a/src/plugins/intel_gpu/src/graph/primitive_inst.cpp b/src/plugins/intel_gpu/src/graph/primitive_inst.cpp index 0b8e4b739f0cf7..d1cb04375a7c49 100644 --- a/src/plugins/intel_gpu/src/graph/primitive_inst.cpp +++ b/src/plugins/intel_gpu/src/graph/primitive_inst.cpp @@ -2411,11 +2411,11 @@ memory::ptr primitive_inst::allocate_output(engine& _engine, if ((_node.is_output() && is_reorder_weights) || (!_node.is_output() && _node.is_type())) reset = false; GPU_DEBUG_LOG << "[" << _node.id() << ": constant]" << std::endl; - return ov::intel_gpu::allocate_memory_evenif_zero_bytes(_engine, layout, alloc_type, reset); + return _engine.allocate_memory(layout, alloc_type, reset); } } else if (!_node.can_share_buffer() || impl_params.can_be_optimized() || _node.is_output()) { GPU_DEBUG_LOG << "[" << _node.id() << ": output]" << std::endl; - return ov::intel_gpu::allocate_memory_evenif_zero_bytes(_engine, layout, alloc_type, reset); + return _engine.allocate_memory(layout, alloc_type, reset); } else { return get_memory_from_pool(_engine, net_id, @@ -2658,6 +2658,18 @@ bool primitive_inst::is_valid_fusion() const { if (fd.is_type()) can_broadcast = ov::PartialShape::broadcast_merge_into(merged_shape, outer_dep_pshape, fd.typed_desc()->broadcast_spec); + // Check if broadcast happens more than single axis. + // Current gemm_tiled_opt kernel FUSED_OP_LOAD macro cannot support broadcast on dynamic dimension. + if (_node->is_type() && can_broadcast == true && merged_shape.rank().get_length() == outer_dep_pshape.rank().get_length()) { + uint8_t broadcast_more_than_single_axis = 0; + for (int64_t i = 0; i < merged_shape.rank().get_length(); i++) { + if (merged_shape.get_shape().at(i) != outer_dep_pshape.get_shape().at(i)) + broadcast_more_than_single_axis++; + } + if (broadcast_more_than_single_axis > 1) + can_broadcast = false; + } + #ifdef ENABLE_ONEDNN_FOR_GPU // WA for OneDNN binary add fusions: we need to broadcast batch dimension to avoid situation with // batch dimension mismatch in OneDNN tensor descriptors as follow: diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bf_tiled.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bf_tiled.cl index 6a5c9e54a8e904..8b351e133d83ad 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bf_tiled.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bf_tiled.cl @@ -17,8 +17,6 @@ // DISPATCH_FSV - output coordinates for each sub-group are calculated from linearized coordinates // DISPATCH_BSV as if they laid in bs_fs_bsv_fsv format, these macros describe fsv and bsv factors; -#define INPUT_LOAD_SIZE 4 - #if FC_KERNEL_DYNAMIC_QUANTIZE KERNEL(quantize_input)( const __global INPUT0_TYPE* input, @@ -28,40 +26,41 @@ KERNEL(quantize_input)( const uint offset = get_global_id(0); const uint input_offset = offset * QUANTIZE_GROUP_SIZE; - const uint quantize_block = QUANTIZE_GROUP_SIZE / 4; - MAKE_VECTOR_TYPE(INPUT0_TYPE, INPUT_LOAD_SIZE) input_0[quantize_block]; - MAKE_VECTOR_TYPE(DQ_TYPE, INPUT_LOAD_SIZE) quantized_value[quantize_block]; + const uint quantize_block = QUANTIZE_GROUP_SIZE / INPUT_LOAD_SIZE; + MAKE_VECTOR_TYPE(INPUT0_TYPE, INPUT_LOAD_SIZE) input_0; + MAKE_VECTOR_TYPE(DQ_TYPE, INPUT_LOAD_SIZE) quantized_value; INPUT0_TYPE max[quantize_block]; unroll_for (uint i = 0 ; i < quantize_block ; ++i) { - input_0[i] = vload4(0, &input[input_offset + i * 4]); - max[i] = fmax(fmax(fabs(input_0[i][0]), fabs(input_0[i][1])), fmax(fabs(input_0[i][2]), fabs(input_0[i][3]))); + input_0 = vload4(0, &input[input_offset + i * 4]); + max[i] = fmax(fmax(fabs(input_0[0]), fabs(input_0[1])), fmax(fabs(input_0[2]), fabs(input_0[3]))); } - INPUT0_TYPE max_value = 0.001; + INPUT0_TYPE max_value = 0.001h; for (uint i = 0 ; i < quantize_block ; i+=8) { INPUT0_TYPE temp = fmax(fmax(fmax(max[i], max[i+1]), fmax(max[i+2], max[i+3])), fmax(fmax(max[i+4], max[i+5]), fmax(max[i+6], max[i+7]))); max_value = fmax(max_value, temp); } - half quan_scale = (half)max_value / 127; + float quan_scale = (float)max_value / 127.f; #if COMPRESSED_WEIGHTS_INT8 int quantized_sum = 0; #endif for (uint i = 0 ; i < quantize_block ; ++i) { - half4 buff = input_0[i] / (half4)quan_scale; - quantized_value[i] = CAT(CAT(convert_, MAKE_VECTOR_TYPE(DQ_TYPE, INPUT_LOAD_SIZE)), _rte)(buff); + input_0 = vload4(0, &input[input_offset + i * 4]); + float4 buff = convert_float4(input_0) / quan_scale; + quantized_value = CAT(CAT(convert_, MAKE_VECTOR_TYPE(DQ_TYPE, INPUT_LOAD_SIZE)), _rte)(buff); #if COMPRESSED_WEIGHTS_INT8 - quantized_sum += quantized_value[i][0] + quantized_value[i][1] + quantized_value[i][2] + quantized_value[i][3]; + quantized_sum += quantized_value[0] + quantized_value[1] + quantized_value[2] + quantized_value[3]; #endif - vstore4(quantized_value[i], 0, &quantized_input[input_offset + i * 4]); + vstore4(quantized_value, 0, &quantized_input[input_offset + i * 4]); } // Pair of quantizing_scale and quantized activation_sum for each group - quan_var[offset * 2] = quan_scale; + quan_var[offset * 2] = convert_half(quan_scale); #if COMPRESSED_WEIGHTS_INT8 - quan_var[(offset * 2) + 1] = CAT(CAT(convert_, INPUT0_TYPE), _rte)(quantized_sum); + quan_var[(offset * 2) + 1] = convert_half(quantized_sum); #endif } #else // !FC_KERNEL_DYNAMIC_QUANTIZE @@ -808,9 +807,6 @@ inline void FUNC(fc_bf_tiled_kernel_default)( // ===================================================================================================================================== } - - - // Dyc Quantize #if USE_SLM && DYNAMIC_QUANTIZE @@ -974,11 +970,38 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( // ===================================================================================================================================== // Main computation loop const uint iterations = MAIN_LOOP_ELEMENTS_COUNT / TILE_IFM_ELEMENTS_SIZE; // TILE_IFM_ELEMENTS_SIZE : (TILE_IFM * SIMD) - // Each sub-group loads 2 Batch - uint idx_sglid = (sglid * TILE_K) % TILE_IFM_ELEMENTS_SIZE; // same index for sglid 0~7 : to tile_k direction - uint batch_sglid = (sglid * TILE_K) / TILE_IFM_ELEMENTS_SIZE; // 0 to 1 : to batch direction - + // Each sub-group loads 2 Batch + const uint idx_sglid = (sglid * TILE_K) % TILE_IFM_ELEMENTS_SIZE; // same index for sglid 0~7 : to tile_k direction + const uint batch_sglid = (sglid * TILE_K) / TILE_IFM_ELEMENTS_SIZE; // 0 to 1 : to batch direction const uint scale_pitch = (TILE_IN_B_PITCH / QUANTIZE_GROUP_SIZE); + + #if PER_TOKEN_SIZE_DYN_QUANTIZE + // Each token is quantized by once. So, all MAIN_LOOP_ELEMENTS_COUNT share just one quantizing variable + uint per_token_offset = input_offset / QUANTIZE_GROUP_SIZE; + unroll_for (uint bi = 0; bi < TILE_B; ++bi) { + de_quantize_scale[bi] = TO_INPUT0_TYPE(quan_var[per_token_offset * 2]); + #if COMPRESSED_WEIGHTS_INT8 + activation_sum[bi] = TO_INPUT0_TYPE(quan_var[per_token_offset * 2 + 1]); + #endif + per_token_offset += scale_pitch; + } + #endif + + #if COMPRESSED_WEIGHTS_INT8 + ACCUMULATOR_TYPE wei_zp[TILE_OFM] = { }; + unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { + #if DECOMPRESSION_ZP_TERM + #if DECOMPRESSION_ZP_SCALAR + wei_zp[fi] = (TO_ACCUMULATOR_TYPE)(DECOMPRESSION_ZP_VALUE); + #elif DECOMPRESSION_ZP_GROUPS_NUM == 1 + wei_zp[fi] = TO_ACCUMULATOR_TYPE(d_zps[fi % DECOMPRESSION_ZP_LENGTH]); + #endif + #else + wei_zp[fi] = ACCUMULATOR_VAL_ZERO; + #endif + } + #endif + MAKE_VECTOR_TYPE(int, TILE_B) acc_tmp[TILE_OFM] = { }; __attribute__((opencl_unroll_hint(1))) for (uint ni = 0; ni < iterations; ++ni) { @@ -993,7 +1016,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( // Next batch in_offset += (TILE_IN_B_PITCH * 2); - #if NUM_LOOP_IN_DYN_QUAN_GROUP == 1 + #if !PER_TOKEN_SIZE_DYN_QUANTIZE && (NUM_LOOP_IN_DYN_QUAN_GROUP == 1) de_quantize_scale[bi * 2] = quan_var[scale_offset * 2]; de_quantize_scale[bi * 2 + 1] = quan_var[scale_offset * 2 + scale_pitch * 2]; #if COMPRESSED_WEIGHTS_INT8 @@ -1006,7 +1029,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif } - #if NUM_LOOP_IN_DYN_QUAN_GROUP > 1 + #if !PER_TOKEN_SIZE_DYN_QUANTIZE && (NUM_LOOP_IN_DYN_QUAN_GROUP > 1) if (ni % NUM_LOOP_IN_DYN_QUAN_GROUP == 0) { unroll_for (uint bi = 0; bi < TILE_B; ++bi) { de_quantize_scale[bi] = quan_var[scale_offset * 2]; @@ -1045,10 +1068,6 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif uint wei_local_idx = local_id * SIMD * FILTER_LOAD_ITERS * (FILTER_LOAD_BLOCK_SIZE/2) + sglid * 2; - #if COMPRESSED_WEIGHTS_INT8 - ACCUMULATOR_TYPE wei_zp[TILE_OFM] = { }; - #endif - // DQ_DECOMPRESSION_SCALE_POST_OP SHOULD be enabled for dynamic quantize FC : scale is ACCUMULATOR_VAL_ONE unroll_for(uint load_iter = 0; load_iter < FILTER_LOAD_ITERS; ++load_iter) { #if COMPRESSED_WEIGHTS_INT4 @@ -1110,31 +1129,6 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif #endif - #if COMPRESSED_WEIGHTS_INT8 - unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { - #if DECOMPRESSION_ZP_TERM - #if DECOMPRESSION_ZP_SCALAR - wei_zp[fi] = (TO_ACCUMULATOR_TYPE)(DECOMPRESSION_ZP_VALUE); - #elif DECOMPRESSION_ZP_GROUPS_NUM > 1 - #if FILTER_LOAD_BLOCK_SIZE % DECOMPRESSION_ZP_GROUP_SIZE != 0 - #error "FC bf_tiled kernel: Not support DECOMPRESSION_ZP_GROUPS_NUM > 1" - #endif - - const uint ni_offset = ni * TILE_IFM * SIMD + local_id * FILTER_LOAD_ITERS * FILTER_LOAD_BLOCK_SIZE; - const uint offset_ofm = out_f + fi*SIMD + sglid; - const uint offset_ifm = ni_offset + load_iter * FILTER_LOAD_BLOCK_SIZE; - const uint zp_offset = (offset_ofm % DECOMPRESSION_ZP_BATCH_NUM) * DECOMPRESSION_ZP_BATCH_PITCH + - (offset_ifm / DECOMPRESSION_ZP_GROUP_SIZE) * DECOMPRESSION_ZP_FEATURE_PITCH; - wei_zp[fi] = TO_ACCUMULATOR_TYPE(decompression_zp[zp_offset]); - #else - wei_zp[fi] = TO_ACCUMULATOR_TYPE(d_zps[fi % DECOMPRESSION_ZP_LENGTH]); - #endif - #else - wei_zp[fi] = ACCUMULATOR_VAL_ZERO; - #endif - } - #endif - #if FILTER_LOAD_BLOCK_SIZE == 2 SLM_WEIGHT_VEC wei_1 = {dq_wei_unpacked.s01, dq_wei_unpacked.s23}; char_slm_weight[wei_local_idx] = as_uint(wei_1); @@ -1162,6 +1156,21 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #else weights_idx += SIMD * FILTER_ACTUAL_LOAD_BLOCK_SIZE; #endif + + #if COMPRESSED_WEIGHTS_INT8 && DECOMPRESSION_ZP_TERM && DECOMPRESSION_ZP_GROUPS_NUM > 1 && !DECOMPRESSION_ZP_SCALAR + unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { + #if FILTER_LOAD_BLOCK_SIZE % DECOMPRESSION_ZP_GROUP_SIZE != 0 + #error "FC bf_tiled kernel: Not support DECOMPRESSION_ZP_GROUPS_NUM > 1" + #endif + + const uint ni_offset = ni * TILE_IFM * SIMD + local_id * FILTER_LOAD_ITERS * FILTER_LOAD_BLOCK_SIZE; + const uint offset_ofm = out_f + fi*SIMD + sglid; + const uint offset_ifm = ni_offset + load_iter * FILTER_LOAD_BLOCK_SIZE; + const uint zp_offset = (offset_ofm % DECOMPRESSION_ZP_BATCH_NUM) * DECOMPRESSION_ZP_BATCH_PITCH + + (offset_ifm / DECOMPRESSION_ZP_GROUP_SIZE) * DECOMPRESSION_ZP_FEATURE_PITCH; + wei_zp[fi] = TO_ACCUMULATOR_TYPE(decompression_zp[zp_offset]); + } + #endif } wei_local_idx = sglid * 2; @@ -1199,7 +1208,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif #if COMPRESSED_WEIGHTS_INT8 - ACCUM_DQ_TYPE modified_calc_buff = ((int *)(&acc_tmp[fi]))[bi] - ((float)(wei_zp[fi]) * (convert_float)(activation_sum[bi])); + ACCUM_DQ_TYPE modified_calc_buff = ((int *)(&acc_tmp[fi]))[bi] - ((float)(wei_zp[fi]) * activation_sum[bi]); ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += (convert_half)(convert_float(modified_calc_buff) * (float)ds * (float)de_quantize_scale[bi]); #else ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi]; @@ -1210,7 +1219,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif } // Whole tile_k elements of each iteration : ki - #if DQ_DECOMPRESSION_SCALE_POST_OP && (TILE_IFM_ELEMENTS_SIZE <= DECOMPRESSION_SCALE_GROUP_SIZE) + #if !PER_TOKEN_SIZE_DYN_QUANTIZE && DQ_DECOMPRESSION_SCALE_POST_OP && (TILE_IFM_ELEMENTS_SIZE <= DECOMPRESSION_SCALE_GROUP_SIZE) // Dynamic-quantizing group size set to same or smaller than scale group size if ((ni % NUM_LOOP_IN_DYN_QUAN_GROUP) == (NUM_LOOP_IN_DYN_QUAN_GROUP - 1)) { const uint ni_offset = ((ni*TILE_IFM*SIMD) / DECOMPRESSION_SCALE_GROUP_SIZE)*DECOMPRESSION_SCALE_FEATURE_PITCH; @@ -1226,7 +1235,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif #if COMPRESSED_WEIGHTS_INT8 - ACCUM_DQ_TYPE modified_calc_buff = ((int *)(&acc_tmp[fi]))[bi] - ((float)(wei_zp[fi]) * (convert_float)(activation_sum[bi])); + ACCUM_DQ_TYPE modified_calc_buff = ((float)((int *)(&acc_tmp[fi]))[bi]) - ((float)(wei_zp[fi]) * activation_sum[bi]); ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += (convert_half)(convert_float(modified_calc_buff) * (float)ds * (float)de_quantize_scale[bi]); #else ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi]; @@ -1238,6 +1247,20 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif } // Main compute loop : ni + #if PER_TOKEN_SIZE_DYN_QUANTIZE + unroll_for (uint bi = 0; bi < TILE_B; ++bi) { + unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { + ACCUMULATOR_TYPE ds = d_scales[fi % DECOMPRESSION_SCALE_LENGTH]; + #if COMPRESSED_WEIGHTS_INT8 + float modified_calc_buff = ((float)((int *)(&acc_tmp[fi]))[bi]) - ((float)(wei_zp[fi]) * activation_sum[bi]); + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] = (convert_half)(modified_calc_buff) * ds * de_quantize_scale[bi]; + #else + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] = convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi]; + #endif + } + } + #endif + // ===================================================================================================================================== // Post-processing: bias, activation, fused-ops for (uint bi = 0; bi < TILE_B; ++bi) { diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.cpp index 0774c62add1643..51dfac1b2b4fee 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.cpp @@ -11,7 +11,8 @@ static constexpr size_t lws_batches = 8; static constexpr size_t simd = 16; -static constexpr size_t min_quantize_grp_size = 32; +static constexpr size_t input_load_size = 4; +static constexpr size_t min_quantize_grp_size = (simd * 2); // SIMD * (min value of tile_ifm) static constexpr size_t min_slm_size = 256; static std::vector available_quantize_grp_size = {128, 64, 32}; @@ -52,17 +53,38 @@ static std::pair get_output_aligned_bf_size(const fully_connecte return {output_b, output_f}; } +static size_t get_scale_group_size(const fully_connected_params& params) { + return params.weights.IFM().v / params.decompression_scale.Feature().v; +} + +static bool is_8bit_asym_wei(const fully_connected_params& params) { + auto weight_type = params.weights.GetDType(); + // UINT8 weight type is supported by FC dyn-quantize(with SLM). + if (weight_type == WeightsType::UINT8 && params.has_decompression_zp) + return true; + + return false; +} + static bool is_weight_dyn_quantizable(const fully_connected_params& params) { auto weight_type = params.weights.GetDType(); if (weight_type == WeightsType::INT4 || weight_type == WeightsType::UINT4) return true; - // UINT8 weight type is supported by FC dyn-quantize(with SLM). - if (weight_type == WeightsType::UINT8) + // No validated case of sym 8bit weight + if (is_8bit_asym_wei(params)) return true; return false; } +static bool is_per_token_dynamic_quantize(const fully_connected_params& params) { + auto dynamic_quantization_group_size = params.dynamic_quantization_group_size; + if (dynamic_quantization_group_size == UINT64_MAX) + return true; + + return false; + } + // DYNAMIC_QUANTIZE static size_t get_dynamic_quantize_group_size(const fully_connected_params& params) { auto dynamic_quantization_group_size = params.dynamic_quantization_group_size; @@ -87,9 +109,35 @@ static size_t get_dynamic_quantize_group_size(const fully_connected_params& para } } - const size_t scale_group_size = params.weights.IFM().v / params.decompression_scale.Feature().v; + size_t scale_group_size = get_scale_group_size(params); + size_t zp_group_num = params.decompression_zero_point.Feature().v; + size_t zp_group_size = 0; + if (params.has_decompression_zp) + zp_group_size = params.weights.IFM().v / params.decompression_zero_point.Feature().v; + + // Per-token dyn-quan + if (dynamic_quantization_group_size >= min_quantize_grp_size && is_per_token_dynamic_quantize(params)) { + // Validate size to fit dyn-quan group to the size of weight-scale and weight-zp + if ((scale_group_size % min_quantize_grp_size) == 0 && scale_group_size > min_quantize_grp_size) { + dynamic_quantization_group_size = scale_group_size; + + // For int8 ASYM model, activation_sum should fit to weight zp + if (is_8bit_asym_wei(params) && params.has_decompression_zp == true && + dynamic_quantization_group_size > zp_group_size && (zp_group_size % input_load_size) == 0) { + dynamic_quantization_group_size = zp_group_size; + } + + GPU_DEBUG_LOG << "FC dyn-quantize by per-token. Actual dyn_quan_group_size(" << dynamic_quantization_group_size + << ") : From scale_group_size (" << scale_group_size << ", zp_group_size(" << zp_group_size + << "), zp_group_num(" << zp_group_num << "), ifm_size (" << get_input_bf_size(params).second << ")" << std::endl; + return (size_t)dynamic_quantization_group_size; + } + } + + // Grouped-size dyn-quan : use aligned sizes which are in 'available_quantize_grp_size' for (auto group_size : available_quantize_grp_size) { - if (dynamic_quantization_group_size >= group_size) { + if (dynamic_quantization_group_size >= group_size && + (scale_group_size % group_size) == 0) { dynamic_quantization_group_size = group_size; if (dynamic_quantization_group_size > scale_group_size) { @@ -97,6 +145,7 @@ static size_t get_dynamic_quantize_group_size(const fully_connected_params& para << dynamic_quantization_group_size << ". Reduce FC dyn-quan group size to scale size." << std::endl; dynamic_quantization_group_size = scale_group_size; } + return (size_t)dynamic_quantization_group_size; } } @@ -104,7 +153,7 @@ static size_t get_dynamic_quantize_group_size(const fully_connected_params& para return 0; } -static bool should_dynamic_quantize(const fully_connected_params& params, bool print_log = false) { +static bool should_dynamic_quantize(const fully_connected_params& params) { size_t dynamic_quantization_group_size = get_dynamic_quantize_group_size(params); if (params.inputs[0].GetFirstElementOffset() != 0) @@ -116,24 +165,25 @@ static bool should_dynamic_quantize(const fully_connected_params& params, bool p return false; } + const size_t scale_group_size = get_scale_group_size(params); + if ((scale_group_size % min_quantize_grp_size) != 0) + return false; + auto threads = get_input_bf_size(params); auto input_b = threads.first; auto input_f = threads.second; - const size_t scale_group_size = params.weights.IFM().v / params.decompression_scale.Feature().v; if ((scale_group_size % simd == 0) && (input_f % dynamic_quantization_group_size == 0) && (params.is_shape_agnostic || (params.inputs[0].Batch().v > 1 && input_b > min_slm_size)) && params.inputs[0].GetDType() == Datatype::F16 && is_weight_dyn_quantizable(params)) { - if (print_log) { - GPU_DEBUG_TRACE_DETAIL << " Dynamic quantizing for FC : scale_group_size: " << scale_group_size << - ", Dyn-quan group size: " << dynamic_quantization_group_size << - ", Type(I:" << kernel_selector::toString(params.inputs[0].GetDType()) << - ", O:" << kernel_selector::toString(params.outputs[0].GetDType()) << - ", W:" << kernel_selector::toString(params.weights.GetDType()) << - "), Format(W:" << kernel_selector::toString(params.weights.GetLayout()) << - ") B: " << params.inputs[0].Batch().v << ", F: " << params.inputs[0].Feature().v << - ", Y: " << params.inputs[0].Y().v << std ::endl; - } + GPU_DEBUG_TRACE_DETAIL << " Dynamic quantizing for FC : scale_group_size: " << scale_group_size << + ", Dyn-quan group size: " << dynamic_quantization_group_size << + ", Type(I:" << kernel_selector::toString(params.inputs[0].GetDType()) << + ", O:" << kernel_selector::toString(params.outputs[0].GetDType()) << + ", W:" << kernel_selector::toString(params.weights.GetDType()) << + "), Format(W:" << kernel_selector::toString(params.weights.GetLayout()) << + ") B: " << params.inputs[0].Batch().v << ", F: " << params.inputs[0].Feature().v << + ", Y: " << params.inputs[0].Y().v << std ::endl; return true; } @@ -621,7 +671,7 @@ JitConstants FullyConnected_bf_tiled::GetJitConstants(const fully_connected_para if (weights_dt == WeightsType::UINT4 || weights_dt == WeightsType::INT4) { tile_k_ofm_packed /= 2; jit.Merge(make_int4_packed_type_jit_constant("INT4_PACKED_TYPE", weights_dt, tile_k_ofm)); - const size_t scale_group_size = params.weights.IFM().v / params.decompression_scale.Feature().v; + const size_t scale_group_size = get_scale_group_size(params); // Do not use SCALE_POST_OP for SLM kernel, since it demonstrates worse performance if (scale_group_size % simd == 0 && !dispatchData.use_slm) add_decompress_scale_post_op = true; @@ -703,16 +753,21 @@ JitConstants FullyConnected_bf_tiled::GetJitConstants(const fully_connected_para jit.AddConstant(MakeJitConstant("DYNAMIC_QUANTIZE", 1)); jit.AddConstant(MakeJitConstant("DQ_DECOMPRESSION_SCALE_POST_OP", 1)); jit.AddConstant(MakeJitConstant("QUANTIZE_GROUP_SIZE", quantize_grp_size)); + + jit.AddConstant(MakeJitConstant("PER_TOKEN_SIZE_DYN_QUANTIZE", + is_per_token_dynamic_quantize(params) && quantize_grp_size == get_input_bf_size(params).second)); } else { if (add_decompress_scale_post_op) jit.AddConstant(MakeJitConstant("DECOMPRESSION_SCALE_POST_OP", 1)); jit.AddConstant(MakeJitConstant("DYNAMIC_QUANTIZE", 0)); jit.AddConstant(MakeJitConstant("QUANTIZE_GROUP_SIZE", min_quantize_grp_size)); } - jit.AddConstant(MakeJitConstant("DQ_TYPE", "char")); + jit.AddConstant(MakeJitConstant("INPUT_LOAD_SIZE", input_load_size)); + jit.AddConstant(MakeJitConstant("DQ_TYPE", "char")); jit.AddConstant(MakeJitConstant("IFM_SIZE", get_input_bf_size(params).second)); jit.AddConstant(MakeJitConstant("SIMD", simd)); + jit.AddConstant(MakeJitConstant("TILE_B", dispatchData.tile_m)); jit.AddConstant(MakeJitConstant("HALF_TILE_B", dispatchData.tile_m/2)); jit.AddConstant(MakeJitConstant("TILE_OFM", dispatchData.tile_n)); @@ -840,19 +895,21 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { kd.kernels[0].skip_execution = false; size_t input_f = get_input_bf_size(prim_params).second; size_t input_size = input_f * dispatchData.tile_m * dispatchData.gws[2]; + OPENVINO_ASSERT(quantize_grp_size != 0, "Error: quantize_grp_size is zero."); + // half type of de_quan_scale and activation sum for each quantized group + size_t quan_var_size = (input_size / quantize_grp_size) * 2 * 2; - if (kd.internalBufferSizes[0] < input_size) { + if (kd.internalBufferSizes[0] < input_size || + kd.internalBufferSizes[1] < quan_var_size) { kd.internalBufferSizes.clear(); // quantized input is char type kd.internalBufferSizes.push_back(input_size); - // half type of de_quan_scale and activation sum for each quantized group - OPENVINO_ASSERT(quantize_grp_size != 0, "Error: quantize_grp_size is zero."); - kd.internalBufferSizes.push_back((input_size / quantize_grp_size) * 2 * 2); + // float type of de_quan_scale and activation sum for each quantized group + kd.internalBufferSizes.push_back(quan_var_size); } - OPENVINO_ASSERT(quantize_grp_size != 0, "Error: quantize_grp_size is zero."); - kd.kernels[0].params.workGroups.global = {std::max((input_size / quantize_grp_size), (size_t)1), 1, 1}; - kd.kernels[0].params.workGroups.local = {16, 1, 1}; + kd.kernels[0].params.workGroups.global = {(std::max((input_size / quantize_grp_size), (size_t)1)), 1, 1}; + kd.kernels[0].params.workGroups.local = {1, 1, 1}; } } }; @@ -1027,8 +1084,10 @@ KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params ¶ms, auto input_size = std::max(fc_params.inputs[0].PhysicalSize(), get_input_bf_size(fc_params).second); if (!params.is_shape_agnostic) input_size = std::max(input_size, Align(get_input_bf_size(fc_params).first, lws_batches) * get_input_bf_size(fc_params).second); - dyn_quan_dispatch.gws = {input_size / quantize_grp_size, 1, 1}; - dyn_quan_dispatch.lws = {16, 1, 1}; + + dyn_quan_dispatch.gws = {(input_size / quantize_grp_size), 1, 1}; + dyn_quan_dispatch.lws = {1, 1, 1}; + quan_kernel.params.workGroups.global = dyn_quan_dispatch.gws; quan_kernel.params.workGroups.local = dyn_quan_dispatch.lws; quan_kernel.skip_execution = false; @@ -1059,7 +1118,7 @@ KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params ¶ms, // char type quantized input kd.internalBufferSizes.push_back(input_size); // half type of de_quan_scale and activation sum for each quantized group - kd.internalBufferSizes.push_back(input_size / quantize_grp_size * 2 * 2); + kd.internalBufferSizes.push_back((input_size / quantize_grp_size) * 2 * 2); kernel_number++; } kd.internalBufferDataType = Datatype::F16; diff --git a/src/plugins/intel_gpu/src/plugin/transformations/clamp_fp16_output.cpp b/src/plugins/intel_gpu/src/plugin/transformations/clamp_fp16_output.cpp index 22def7242865a7..ba779eb0f52ca9 100644 --- a/src/plugins/intel_gpu/src/plugin/transformations/clamp_fp16_output.cpp +++ b/src/plugins/intel_gpu/src/plugin/transformations/clamp_fp16_output.cpp @@ -31,11 +31,11 @@ ClampFP16Output::ClampFP16Output() { auto in0 = any_input(as_value_predicate(class_other_than())); auto in1 = any_input(as_value_predicate(class_other_than())); - auto matmul_m = wrap_type({in0, in1}, all_of({type_matches(ov::element::f16), consumers_count(1)})); - auto reshape_m = wrap_type({matmul_m, any_input()}, all_of({type_matches(ov::element::f16), consumers_count(1)})); - auto add_m = wrap_type({matmul_m, any_input()}, all_of({type_matches(ov::element::f16), consumers_count(1)})); + auto matmul_m = wrap_type({in0, in1}, type_matches(ov::element::f16) && consumers_count(1)); + auto reshape_m = wrap_type({matmul_m, any_input()}, type_matches(ov::element::f16) && consumers_count(1)); + auto add_m = wrap_type({matmul_m, any_input()}, type_matches(ov::element::f16) && consumers_count(1)); auto eltwise_m = wrap_type({matmul_m, any_input()}, - all_of({type_matches(ov::element::f16), consumers_count(1)})); + type_matches(ov::element::f16) && consumers_count(1)); auto softmax_input_m = std::make_shared(ov::OutputVector{eltwise_m, reshape_m, matmul_m}); auto softmax_m = wrap_type({softmax_input_m}, type_matches(ov::element::f16)); diff --git a/src/plugins/intel_gpu/src/plugin/transformations/convert_convolution.cpp b/src/plugins/intel_gpu/src/plugin/transformations/convert_convolution.cpp index 0d23eb49829217..32a92c08470d97 100644 --- a/src/plugins/intel_gpu/src/plugin/transformations/convert_convolution.cpp +++ b/src/plugins/intel_gpu/src/plugin/transformations/convert_convolution.cpp @@ -114,11 +114,11 @@ class AsymmetricConvolutionMatcher : public ov::pass::MatcherPass { AsymmetricConvolutionMatcher::AsymmetricConvolutionMatcher() { auto input_m = any_input(type_matches_any({ov::element::u8, ov::element::i8})); - auto azp_const_m = wrap_type(all_of({consumers_count(1), type_matches_any({ov::element::u8, ov::element::i8})})); + auto azp_const_m = wrap_type(consumers_count(1) && type_matches_any({ov::element::u8, ov::element::i8})); auto azp_subtract_m = wrap_type({input_m, azp_const_m}); auto weights_m = wrap_type(type_matches_any({ov::element::u8, ov::element::i8})); - auto wzp_const_m = wrap_type(all_of({consumers_count(1), type_matches_any({ov::element::u8, ov::element::i8})})); + auto wzp_const_m = wrap_type(consumers_count(1) && type_matches_any({ov::element::u8, ov::element::i8})); auto wzp_subtract_m = wrap_type({weights_m, wzp_const_m}); auto conv_activations_m = std::make_shared(OutputVector{input_m, azp_subtract_m}); diff --git a/src/plugins/intel_gpu/src/plugin/transformations/optimize_subsequent_reshapes.cpp b/src/plugins/intel_gpu/src/plugin/transformations/optimize_subsequent_reshapes.cpp index b87600ed36e347..cac1e53761bc1c 100644 --- a/src/plugins/intel_gpu/src/plugin/transformations/optimize_subsequent_reshapes.cpp +++ b/src/plugins/intel_gpu/src/plugin/transformations/optimize_subsequent_reshapes.cpp @@ -41,7 +41,7 @@ OptimizeSubsequentReshapes::OptimizeSubsequentReshapes() { auto first_reshape_data = any_input(dynamic_batch_only); auto first_reshape_pattern = ov::pass::pattern::wrap_type(); auto first_reshape = wrap_type({ first_reshape_data, first_reshape_pattern }, - ov::pass::pattern::all_of({ dynamic_batch_only, ov::pass::pattern::consumers_count(1) })); + dynamic_batch_only && ov::pass::pattern::consumers_count(1)); auto second_reshape_pattern = ov::pass::pattern::wrap_type(); auto second_reshape = wrap_type({ first_reshape, second_reshape_pattern }, dynamic_batch_only); diff --git a/src/plugins/intel_gpu/src/runtime/execution_config.cpp b/src/plugins/intel_gpu/src/runtime/execution_config.cpp index a698ec7eb6c5a0..3b1376d19b4fea 100644 --- a/src/plugins/intel_gpu/src/runtime/execution_config.cpp +++ b/src/plugins/intel_gpu/src/runtime/execution_config.cpp @@ -229,7 +229,27 @@ void ExecutionConfig::apply_hints(const cldnn::device_info& info) { apply_debug_options(info); } +void ExecutionConfig::update_specific_default_properties(const cldnn::device_info& info) { + // These default properties should be set once. + if (specific_default_properties_is_set) + return; + specific_default_properties_is_set = true; + + // Enable KV-cache compression by default for non-systolic platforms + if (get_property(ov::hint::kv_cache_precision) == ov::element::undefined && !info.supports_immad) { + set_property(ov::hint::kv_cache_precision(ov::element::i8)); + } + + // Enable dynamic quantization by default for non-systolic platforms + if (get_property(ov::hint::dynamic_quantization_group_size) == 0 && !info.supports_immad) { + set_property(ov::hint::dynamic_quantization_group_size(32)); + } +} + void ExecutionConfig::apply_user_properties(const cldnn::device_info& info) { + // Update specific default properties, call once before internal_properties updated. + update_specific_default_properties(info); + // Copy internal properties before applying hints to ensure that // a property set by hint won't be overriden by a value in user config. // E.g num_streams=AUTO && hint=THROUGHPUT @@ -249,16 +269,6 @@ void ExecutionConfig::apply_user_properties(const cldnn::device_info& info) { set_property(ov::intel_gpu::queue_type(QueueTypes::in_order)); } - // Enable KV-cache compression by default for non-systolic platforms - if (!is_set_by_user(ov::hint::kv_cache_precision) && !info.supports_immad) { - set_property(ov::hint::kv_cache_precision(ov::element::i8)); - } - - // Enable dynamic quantization by default for non-systolic platforms - if (!is_set_by_user(ov::hint::dynamic_quantization_group_size) && !info.supports_immad) { - set_property(ov::hint::dynamic_quantization_group_size(32)); - } - user_properties.clear(); } diff --git a/src/plugins/intel_gpu/src/runtime/ocl/ocl_memory.cpp b/src/plugins/intel_gpu/src/runtime/ocl/ocl_memory.cpp index ee7e9a85d735b8..0c8fb2126a8aaa 100644 --- a/src/plugins/intel_gpu/src/runtime/ocl/ocl_memory.cpp +++ b/src/plugins/intel_gpu/src/runtime/ocl/ocl_memory.cpp @@ -454,22 +454,25 @@ gpu_usm::gpu_usm(ocl_engine* engine, const layout& layout, allocation_type type) , memory(engine, layout, type, nullptr) , _buffer(engine->get_usm_helper()) , _host_buffer(engine->get_usm_helper()) { + auto actual_bytes_count = _bytes_count; + if (actual_bytes_count == 0) + actual_bytes_count = 1; switch (get_allocation_type()) { case allocation_type::usm_host: - _buffer.allocateHost(_bytes_count); + _buffer.allocateHost(actual_bytes_count); break; case allocation_type::usm_shared: - _buffer.allocateShared(_bytes_count); + _buffer.allocateShared(actual_bytes_count); break; case allocation_type::usm_device: - _buffer.allocateDevice(_bytes_count); + _buffer.allocateDevice(actual_bytes_count); break; default: CLDNN_ERROR_MESSAGE("gpu_usm allocation type", "Unknown unified shared memory type!"); } - m_mem_tracker = std::make_shared(engine, _buffer.get(), layout.bytes_count(), type); + m_mem_tracker = std::make_shared(engine, _buffer.get(), actual_bytes_count, type); } void* gpu_usm::lock(const stream& stream, mem_lock_type type) { diff --git a/src/plugins/intel_gpu/tests/functional/subgraph_tests/dynamic/dynamic_unfusion.cpp b/src/plugins/intel_gpu/tests/functional/subgraph_tests/dynamic/dynamic_unfusion.cpp new file mode 100644 index 00000000000000..1cc079a10b82f6 --- /dev/null +++ b/src/plugins/intel_gpu/tests/functional/subgraph_tests/dynamic/dynamic_unfusion.cpp @@ -0,0 +1,95 @@ +// Copyright (C) 2024 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "common_test_utils/ov_tensor_utils.hpp" +#include "common_test_utils/file_utils.hpp" +#include "shared_test_classes/base/ov_subgraph.hpp" + +#include "openvino/op/parameter.hpp" +#include "openvino/op/matmul.hpp" +#include "openvino/op/multiply.hpp" + +namespace { + +using ov::test::InputShape; + +using DynamicUnfusionsParams = std::tuple, // input shapes + ov::element::Type>; // input precision + +class DynamicUnfusions : public testing::WithParamInterface, + virtual public ov::test::SubgraphBaseTest { +public: + static std::string getTestCaseName(testing::TestParamInfo obj) { + std::vector input_shapes; + ov::element::Type input_precision; + + std::tie(input_shapes, input_precision) = obj.param; + + std::ostringstream result; + result << "IS=("; + for (const auto& shape : input_shapes) { + result << ov::test::utils::partialShape2str({shape.first}) << "_"; + } + result << ")_TS="; + for (const auto& shape : input_shapes) { + result << "("; + if (!shape.second.empty()) { + auto itr = shape.second.begin(); + do { + result << ov::test::utils::vec2str(*itr); + } while (++itr != shape.second.end() && result << "_"); + } + result << ")_"; + } + result << "input_precision=" << input_precision; + return result.str(); + } + +protected: + std::shared_ptr init_subgraph(std::vector& input_shapes, + const ov::element::Type input_precision) { + auto input0 = std::make_shared(input_precision, input_shapes[0]); + auto input1 = std::make_shared(input_precision, input_shapes[1]); + auto input2 = std::make_shared(input_precision, input_shapes[2]); + + auto matmul = std::make_shared(input0, input1); + auto mul = std::make_shared(matmul, input2); + + matmul->set_friendly_name("MatMul"); + mul->set_friendly_name("Multiply"); + + return std::make_shared(ov::NodeVector{mul}, ov::ParameterVector{input0, input1, input2}, "DynamicUnfusions"); + } + + void SetUp() override { + targetDevice = ov::test::utils::DEVICE_GPU; + + std::vector input_shapes; + ov::element::Type input_precision; + + std::tie(input_shapes, input_precision) = GetParam(); + + init_input_shapes(input_shapes); + + inType = outType = input_precision; + function = init_subgraph(inputDynamicShapes, input_precision); + } +}; + +TEST_P(DynamicUnfusions, Inference) { + run(); +} + +const std::vector input_precisions = {ov::element::f32}; + +const std::vector> input_shapes_dyn = { + {{{1024, -1}, {{1024, 1024}}}, {{-1, 1024}, {{1024, 1024}}}, {{1, -1}, {{1, 1}}}}, +}; + +INSTANTIATE_TEST_SUITE_P(DynamicUnfusions_basic, + DynamicUnfusions, + ::testing::Combine(::testing::ValuesIn(input_shapes_dyn), + ::testing::ValuesIn(input_precisions)), + DynamicUnfusions::getTestCaseName); +} // namespace diff --git a/src/plugins/intel_gpu/tests/unit/test_cases/fully_connected_gpu_test.cpp b/src/plugins/intel_gpu/tests/unit/test_cases/fully_connected_gpu_test.cpp index 5bc7e403d3bf74..8a90d137ed7e5a 100644 --- a/src/plugins/intel_gpu/tests/unit/test_cases/fully_connected_gpu_test.cpp +++ b/src/plugins/intel_gpu/tests/unit/test_cases/fully_connected_gpu_test.cpp @@ -2846,7 +2846,7 @@ class fully_connected_gpu_tests: public ::testing::Test { } void test_compressed_int4_scale_dyn_quan_weight_i4(bool is_dynamic, int batch = 1, int ifm = 512, int ofm = 2048, - int quantize_group_size = 32, int scales_group_size = 128, + size_t quantize_group_size = 32, int scales_group_size = 128, bool is_wzp_test = false, bool is_wzp_scalar = false) { tests::random_generator rg(GET_SUITE_NAME); auto& engine = get_test_engine(); @@ -2972,7 +2972,7 @@ class fully_connected_gpu_tests: public ::testing::Test { } void test_compressed_int8_scale_dyn_quan_weight_u8(bool is_dynamic, int batch = 1, int ifm = 512, int ofm = 2048, - int quantize_group_size = 32, int scales_group_size = 128, + size_t quantize_group_size = 32, int scales_group_size = 128, bool is_wzp_test = false, bool is_wzp_scalar = false) { tests::random_generator rg(GET_SUITE_NAME); auto& engine = get_test_engine(); @@ -3065,8 +3065,9 @@ class fully_connected_gpu_tests: public ::testing::Test { auto inst = network->get_primitive("fc_prim"); auto impl = inst->get_impl(); ASSERT_TRUE(impl != NULL); - auto kernel_num = (is_dynamic) ? 3 : 2; - kernel_num = (quantize_group_size < 32) ? 2 : kernel_num; + // For UINT8 weight, SLM kernel (no dyn-quan) would not be selected + auto kernel_num = (is_dynamic) ? 3 : 1; + kernel_num = (quantize_group_size < 32) ? 1 : kernel_num; ASSERT_EQ(impl->get_kernels().size(), size_t(kernel_num)); } @@ -4194,6 +4195,7 @@ TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dynamic_quantize_wzp_sta this->test_compressed_int4_scale_dyn_quan_weight_i4(false, 320, 1024, 1024, 32, 32, true); } +// Test weight zp for INT8 ASYM TEST_F(fully_connected_gpu_tests, compressed_int8_scale_dynamic_quantize_wzp_128_large) { this->test_compressed_int8_scale_dyn_quan_weight_u8(true, 320, 4096, 4096, 128, 128, true); } @@ -4210,10 +4212,6 @@ TEST_F(fully_connected_gpu_tests, compressed_int8_scale_dynamic_quantize_wzp_32_ this->test_compressed_int8_scale_dyn_quan_weight_u8(true, 320, 4096, 4096, 32, 32, true); } -TEST_F(fully_connected_gpu_tests, compressed_int8_scale_dynamic_quantize_wzp_32_large_unaligned) { - this->test_compressed_int8_scale_dyn_quan_weight_u8(true, 310, 1024, 1024, 32, 32, true); -} - TEST_F(fully_connected_gpu_tests, compressed_int8_scale_dynamic_quantize_wzp_128_small) { this->test_compressed_int8_scale_dyn_quan_weight_u8(true, 16, 1024, 1024, 128, 128, true); } @@ -4222,6 +4220,23 @@ TEST_F(fully_connected_gpu_tests, compressed_int8_scale_dynamic_quantize_wzp_128 this->test_compressed_int8_scale_dyn_quan_weight_u8(true, 1, 1024, 1024, 128, 128, true); } +// Test per-token dyn-quan +TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dynamic_quantize_test_fake_per_token) { + this->test_compressed_int4_scale_dyn_quan_weight_i4(true, 600, 1024, 2048, -1, 32, true); +} + +TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dynamic_quantize_test_per_token) { + this->test_compressed_int4_scale_dyn_quan_weight_i4(true, 600, 1024, 2048, -1, 1024, true); +} + +TEST_F(fully_connected_gpu_tests, compressed_int8_scale_dynamic_quantize_test_per_token_small_scale) { + this->test_compressed_int8_scale_dyn_quan_weight_u8(true, 600, 1024, 2048, -1, 32, true); +} + +TEST_F(fully_connected_gpu_tests, compressed_int8_scale_dynamic_quantize_test_per_token_full_scale) { + this->test_compressed_int8_scale_dyn_quan_weight_u8(true, 600, 1024, 2048, -1, 1024, true); +} + TEST_F(fully_connected_gpu_tests, compressed_scale_bias) { this->test_compressed_scale_bias(false); } diff --git a/src/plugins/intel_npu/src/al/include/intel_npu/config/npuw.hpp b/src/plugins/intel_npu/src/al/include/intel_npu/config/npuw.hpp index 927b234df8ba15..51b6b8ca9fe28c 100644 --- a/src/plugins/intel_npu/src/al/include/intel_npu/config/npuw.hpp +++ b/src/plugins/intel_npu/src/al/include/intel_npu/config/npuw.hpp @@ -68,60 +68,18 @@ DEFINE_OPT(NPUW_DUMP_SUBS_ON_FAIL, std::string, "", npuw::dump::subgraphs_on_fai DEFINE_OPT(NPUW_DUMP_IO, std::string, "", npuw::dump::inputs_outputs, RunTime); DEFINE_OPT(NPUW_DUMP_IO_ITERS, bool, false, npuw::dump::io_iters, RunTime); DEFINE_OPT(NPUW_LLM, bool, false, npuw::llm::enabled, CompileTime); +DEFINE_OPT(NPUW_LLM_BATCH_DIM, uint32_t, 0, npuw::llm::batch_dim, CompileTime); +DEFINE_OPT(NPUW_LLM_SEQ_LEN_DIM, uint32_t, 2, npuw::llm::seq_len_dim, CompileTime); DEFINE_OPT(NPUW_LLM_MAX_PROMPT_LEN, uint32_t, 1024, npuw::llm::max_prompt_len, CompileTime); DEFINE_OPT(NPUW_LLM_MIN_RESPONSE_LEN, uint32_t, 128, npuw::llm::min_response_len, CompileTime); +DEFINE_OPT(NPUW_LLM_OPTIMIZE_V_TENSORS, bool, false, npuw::llm::optimize_v_tensors, CompileTime); namespace npuw { namespace llm { -struct ModelDesc { - std::string type; - std::string name_or_path; - int num_key_value_heads; -}; enum class GenerateHint { FAST_COMPILE, BEST_PERF }; } // namespace llm } // namespace npuw -struct NPUW_LLM_MODEL_DESC final : OptionBase { - static std::string_view key() { - return ov::intel_npu::npuw::llm::model_desc.name(); - } - - static constexpr std::string_view getTypeName() { - return "::intel_npu::npuw::llm::ModelDesc"; - } - - static ::intel_npu::npuw::llm::ModelDesc defaultValue() { - return {}; - } - - static ::intel_npu::npuw::llm::ModelDesc parse(std::string_view val) { - ::intel_npu::npuw::llm::ModelDesc res; - std::map res_map = OptionParser>::parse(val); - res.type = res_map["type"]; - res.name_or_path = res_map["name_or_path"]; - res.num_key_value_heads = std::stoi(res_map["num_key_value_heads"]); - return res; - } - - static std::string toString(const ::intel_npu::npuw::llm::ModelDesc& val) { - std::string res; - std::map res_map; - res_map["type"] = val.type; - res_map["name_or_path"] = val.name_or_path; - res_map["num_key_value_heads"] = std::to_string(val.num_key_value_heads); - return OptionPrinter>::toString(res_map); - } - - static OptionMode mode() { - return OptionMode::CompileTime; - } - - static bool isPublic() { - return true; - } -}; - struct NPUW_LLM_GENERATE_HINT final : OptionBase { static std::string_view key() { return ov::intel_npu::npuw::llm::generate_hint.name(); diff --git a/src/plugins/intel_npu/src/al/include/intel_npu/npuw_private_properties.hpp b/src/plugins/intel_npu/src/al/include/intel_npu/npuw_private_properties.hpp index a416ca51233893..5cf489b6df34b4 100644 --- a/src/plugins/intel_npu/src/al/include/intel_npu/npuw_private_properties.hpp +++ b/src/plugins/intel_npu/src/al/include/intel_npu/npuw_private_properties.hpp @@ -389,17 +389,26 @@ static constexpr ov::Property enabled{"NPUW_LLM"}; /** * @brief - * Type: std::map. - * Tell NPUW about your LLM model. Use following structure for that: - * "type:,name_or_path:,num_key_value_heads:". - * Default value: empty structure defined above. + * FIXME: Should be removed. + * Type: uint32_t. + * Dimension of the batch in input tensor shape. + * Default value: 0. + */ +static constexpr ov::Property batch_dim{"NPUW_LLM_BATCH_DIM"}; + +/** + * @brief + * FIXME: Should be removed. + * Type: uint32_t. + * Dimension of KV-Cache size in input tensor shape. + * Default value: 2. */ -static constexpr ov::Property model_desc{"NPUW_LLM_MODEL_DESC"}; +static constexpr ov::Property seq_len_dim{"NPUW_LLM_SEQ_LEN_DIM"}; /** * @brief * Type: uint32_t. - * Tell NPUW your desirable max prompt length. + * Desirable max prompt length. * Default value: 1024. */ static constexpr ov::Property max_prompt_len{"NPUW_LLM_MAX_PROMPT_LEN"}; @@ -407,20 +416,45 @@ static constexpr ov::Property max_prompt_len{"NPUW_LLM_MAX_PROMPT_LEN" /** * @brief * Type: uint32_t. - * Tell NPUW your desirable min response length. + * Desirable min response length. * Default value: 128. */ static constexpr ov::Property min_response_len{"NPUW_LLM_MIN_RESPONSE_LEN"}; +/** + * @brief + * FIXME: Should be removed. + * Type: bool. + * Tell NPUW to apply values transpose optimization for the model. + * Default value: false. + */ +static constexpr ov::Property optimize_v_tensors{"NPUW_LLM_OPTIMIZE_V_TENSORS"}; + +/** + * @brief + * Type: ov::AnyMap. + * Configuration for compilation of prefill model. + * NOTE: !! Write-only !! + */ +static constexpr ov::Property prefill_config{"NPUW_LLM_PREFILL_CONFIG"}; + /** * @brief * Type: std::string. - * Tell NPUW the preferrable hint for generation stage, that leads to usage of optimal configuration for it. + * Hint for generation stage. NPUW will use optimal configuration based on the passed preference via hint. + * Hint is ignored if used with "NPUW_LLM_GENERATE_CONFIG". * Possible values: "FAST_COMPILE", "BEST_PERF". * Default value: "FAST_COMPILE". */ static constexpr ov::Property generate_hint{"NPUW_LLM_GENERATE_HINT"}; +/** + * @brief + * Type: ov::AnyMap. + * Configuration for compilation of generate model. + * NOTE: !! Write-only !! + */ +static constexpr ov::Property generate_config{"NPUW_LLM_GENERATE_CONFIG"}; } // namespace llm } // namespace npuw diff --git a/src/plugins/intel_npu/src/al/src/config/npuw.cpp b/src/plugins/intel_npu/src/al/src/config/npuw.cpp index 4ee9e392406452..963afb5a3efe15 100644 --- a/src/plugins/intel_npu/src/al/src/config/npuw.cpp +++ b/src/plugins/intel_npu/src/al/src/config/npuw.cpp @@ -57,8 +57,10 @@ void intel_npu::registerNPUWOptions(OptionsDesc& desc) { void intel_npu::registerNPUWLLMOptions(OptionsDesc& desc) { desc.add(); - desc.add(); + desc.add(); + desc.add(); desc.add(); desc.add(); + desc.add(); desc.add(); } 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 eb13bc8b5bd1d9..cea6a38041d1fc 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 @@ -280,22 +280,6 @@ void reshape_to_static(std::shared_ptr model, model->reshape(new_shapes); } -KVAxesPosition get_kv_axes(const std::string& model_type) { - KVAxesPosition axes; - if (model_type == "chatglm") { - axes.batch = 1u; - axes.seq_len = 0u; - } else if (model_type == "qwen") { - // Note, qwen2 does not fall into this category and conforms to default layout - axes.batch = 0u; - axes.seq_len = 1u; - } else { - axes.batch = 0u; - axes.seq_len = 2u; - } - return axes; -} - bool is_cw_compressed(const std::shared_ptr& model) { std::vector rt_info_path = {"nncf", "weight_compression", "group_size"}; if (!model->has_rt_info(rt_info_path)) { @@ -321,6 +305,15 @@ 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; +} + ov::AnyMap get_baseline_common_config() { ov::AnyMap config = { {"NPU_COMPILATION_MODE_PARAMS", "compute-layers-with-higher-precision=Sqrt,Power,ReduceMean,Add_RMSNorm"}, @@ -418,6 +411,13 @@ ov::npuw::LLMCompiledModel::LLMCompiledModel(const std::shared_ptr& m std::map npuw_llm_props; std::map other_props; split_llm_properties(properties, npuw_llm_props, other_props); + + // Remove "NPUW_LLM_PREFILL_CONFIG", "NPUW_LLM_GENERATE_CONFIG" from map, + // to not pass them into ::intel_npu::Config object, as we don't need to + // preserve them somewhere. + auto prefill_config_opt = pop_option(npuw_llm_props, std::string("NPUW_LLM_PREFILL_CONFIG")); + auto generate_config_opt = pop_option(npuw_llm_props, std::string("NPUW_LLM_GENERATE_CONFIG")); + m_cfg.update(any_copy(npuw_llm_props)); LOG_DEBUG("1. Creating kvcache model as clone of passed one."); @@ -428,19 +428,22 @@ ov::npuw::LLMCompiledModel::LLMCompiledModel(const std::shared_ptr& m auto prefill_model = kvcache_model->clone(); prefill_model->set_friendly_name(kvcache_model->get_friendly_name() + "_prefill"); - 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); - KVAxesPosition axes = get_kv_axes(model_desc.type); - m_kvcache_desc = KVCacheDesc{kMaxPromptLen, kMaxPromptLen + kMinResponseLen, 0u, axes.seq_len}; + const uint32_t batch_dim = m_cfg.get<::intel_npu::NPUW_LLM_BATCH_DIM>(); + const uint32_t seq_len_dim = m_cfg.get<::intel_npu::NPUW_LLM_SEQ_LEN_DIM>(); + KVAxesPosition axes{batch_dim, seq_len_dim}; + const uint32_t max_prompt_len = align_to(m_cfg.get<::intel_npu::NPUW_LLM_MAX_PROMPT_LEN>(), 64u); + const uint32_t min_response_len = align_to(m_cfg.get<::intel_npu::NPUW_LLM_MIN_RESPONSE_LEN>(), 64u); + + m_kvcache_desc = KVCacheDesc{max_prompt_len, max_prompt_len + min_response_len, 0u, seq_len_dim}; 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("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."); + + const bool optimize_v_tensors = m_cfg.get<::intel_npu::NPUW_LLM_OPTIMIZE_V_TENSORS>(); // 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_v_tensors) { if (optimize_value_tensors(kvcache_model)) { // NB: Check if TransposeValueTensors transformation was applied m_kvcache_desc.v_tensors_transposed = true; @@ -455,17 +458,20 @@ ov::npuw::LLMCompiledModel::LLMCompiledModel(const std::shared_ptr& m prefill_model = cvt_kvcache_to_fp16(prefill_model); auto npudesc = extract_npu_descriptor(plugin); - ov::AnyMap properties_copy = other_props; - auto prefill_config = get_default_prefill_config(model, npudesc); + auto prefill_config = + prefill_config_opt.value_or(get_default_prefill_config(prefill_model, npudesc)).as(); - // 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( - "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); + LOG_DEBUG("9. Passed GENERATE_HINT: " << std::string(::intel_npu::NPUW_LLM_GENERATE_HINT::toString(generate_hint))); + // NB: GENERATE_HINT is only applicable for default generate config! + if (generate_config_opt.has_value() && npuw_llm_props.count(ov::intel_npu::npuw::llm::generate_hint.name())) { + OPENVINO_THROW("GENERATE_HINT is only applicable for default generate config!"); + } + auto generate_config = + generate_config_opt.value_or(get_default_generate_config(model, npudesc, generate_hint)).as(); - merge_config_with(prefill_config, properties_copy); - merge_config_with(generate_config, properties_copy); + merge_config_with(prefill_config, other_props); + merge_config_with(generate_config, other_props); m_kvcache_compiled = std::make_shared(kvcache_model, plugin, generate_config); m_prefill_compiled = std::make_shared(prefill_model, plugin, prefill_config); @@ -488,6 +494,11 @@ void ov::npuw::LLMCompiledModel::set_property(const ov::AnyMap& properties) { ov::Any ov::npuw::LLMCompiledModel::get_property(const std::string& name) const { OPENVINO_SUPPRESS_DEPRECATED_START + if (name == ov::intel_npu::npuw::llm::prefill_config.name() || + name == ov::intel_npu::npuw::llm::generate_config.name()) { + OPENVINO_THROW(name, " is write-only option!"); + } + auto&& configIterator = m_prop_to_opt.find(name); if (configIterator != m_prop_to_opt.cend()) { return std::get<1>(configIterator->second)(m_cfg); @@ -504,7 +515,7 @@ std::shared_ptr ov::npuw::LLMCompiledModel::create_sync_i std::shared_ptr ov::npuw::LLMCompiledModel::create_llm_infer_request() { auto this_sptr = std::static_pointer_cast(shared_from_this()); - return std::make_shared(this_sptr, m_kvcache_desc); + return std::make_shared(this_sptr); } void ov::npuw::LLMCompiledModel::implement_properties() { @@ -518,9 +529,11 @@ void ov::npuw::LLMCompiledModel::implement_properties() { } m_prop_to_opt.insert({BIND(npuw::llm::enabled, NPUW_LLM, get), - BIND(npuw::llm::model_desc, NPUW_LLM_MODEL_DESC, getString), + BIND(npuw::llm::batch_dim, NPUW_LLM_BATCH_DIM, get), + BIND(npuw::llm::batch_dim, NPUW_LLM_SEQ_LEN_DIM, get), BIND(npuw::llm::max_prompt_len, NPUW_LLM_MAX_PROMPT_LEN, get), BIND(npuw::llm::min_response_len, NPUW_LLM_MIN_RESPONSE_LEN, get), + BIND(npuw::llm::optimize_v_tensors, NPUW_LLM_OPTIMIZE_V_TENSORS, get), BIND(npuw::llm::generate_hint, NPUW_LLM_GENERATE_HINT, getString)}); #undef BIND } 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 12f103cc0ab6a2..0bad68a35aa4ba 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 @@ -59,10 +59,9 @@ void copy_columns_by_row_chunks(ov::SoPtr src, ov::SoPtr& compiled_model, - const ov::npuw::LLMCompiledModel::KVCacheDesc& kvcache_desc) +ov::npuw::LLMInferRequest::LLMInferRequest(const std::shared_ptr& compiled_model) : ov::ISyncInferRequest(compiled_model), - m_kvcache_desc(kvcache_desc) { + m_npuw_llm_compiled_model(compiled_model) { m_kvcache_request = compiled_model->m_kvcache_compiled->create_infer_request(); m_prefill_request = compiled_model->m_prefill_compiled->create_infer_request(); @@ -82,13 +81,11 @@ ov::npuw::LLMInferRequest::LLMInferRequest(const std::shared_ptr(m_prefill_request->get_tensor(m_prefill_in_ports.at("input_ids")), 0u); - fill_tensor(m_prefill_request->get_tensor(m_prefill_in_ports.at("attention_mask")), 0u); - fill_tensor(m_prefill_request->get_tensor(m_prefill_in_ports.at("position_ids")), 0u); - fill_tensor(m_kvcache_request->get_tensor(m_kvcache_in_ports.at("attention_mask")), 0u); - m_kvcache_desc.num_stored_tokens = 0u; + fill_tensor(m_prefill_request->get_tensor(m_prefill_in_ports.at("input_ids")), 0); + fill_tensor(m_prefill_request->get_tensor(m_prefill_in_ports.at("attention_mask")), 0); + fill_tensor(m_prefill_request->get_tensor(m_prefill_in_ports.at("position_ids")), 0); + fill_tensor(m_kvcache_request->get_tensor(m_kvcache_in_ports.at("attention_mask")), 0); + m_npuw_llm_compiled_model->m_kvcache_desc.num_stored_tokens = 0u; } void ov::npuw::LLMInferRequest::infer_prefill(ov::SoPtr input_ids, @@ -112,7 +109,7 @@ void ov::npuw::LLMInferRequest::infer_prefill(ov::SoPtr input_ids, std::copy_n(position_ids->data(), position_ids->get_size(), padded_position_ids->data() + offset); m_prefill_request->infer(); - m_kvcache_desc.num_stored_tokens += static_cast(input_ids->get_size()); + m_npuw_llm_compiled_model->m_kvcache_desc.num_stored_tokens += static_cast(input_ids->get_size()); m_need_copy_kvcache = true; m_logits = m_prefill_request->get_tensor(m_prefill_out_ports.at("logits")); @@ -126,8 +123,9 @@ void ov::npuw::LLMInferRequest::infer_generate(ov::SoPtr input_ids, LOG_DEBUG("Calling inference for generate model..."); LOG_BLOCK(); + auto& kvcache_desc = m_npuw_llm_compiled_model->m_kvcache_desc; // NB: KV-cache is full, further generation is impossible - if (m_kvcache_desc.num_stored_tokens == m_kvcache_desc.total_size) { + if (kvcache_desc.num_stored_tokens == kvcache_desc.total_size) { OPENVINO_THROW("KV-Cache is full."); } @@ -146,17 +144,16 @@ 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) + const auto& kv_dim = (output_name.find("value") != std::string::npos && kvcache_desc.v_tensors_transposed) ? 3u - : m_kvcache_desc.dim; + : kvcache_desc.dim; - auto prefill_out_slice = - make_tensor_slice(prefill_out_tensor, - kv_dim, - m_kvcache_desc.max_prompt_size - m_kvcache_desc.num_stored_tokens, - m_kvcache_desc.max_prompt_size); + auto prefill_out_slice = make_tensor_slice(prefill_out_tensor, + kv_dim, + kvcache_desc.max_prompt_size - kvcache_desc.num_stored_tokens, + kvcache_desc.max_prompt_size); - auto kvcache_in_slice = make_tensor_slice(kvcache_in_tensor, kv_dim, 0u, m_kvcache_desc.num_stored_tokens); + auto kvcache_in_slice = make_tensor_slice(kvcache_in_tensor, kv_dim, 0u, kvcache_desc.num_stored_tokens); if (kv_dim == 3u) { copy_columns_by_row_chunks(prefill_out_slice, kvcache_in_slice); @@ -168,7 +165,7 @@ void ov::npuw::LLMInferRequest::infer_generate(ov::SoPtr input_ids, LOG_DEBUG("Prepare attention mask pattern."); auto* attention_mask_data = m_kvcache_request->get_tensor(m_kvcache_in_ports.at("attention_mask"))->data(); - attention_mask_data[m_kvcache_desc.total_size - 1] = 1; + attention_mask_data[kvcache_desc.total_size - 1] = 1; m_need_copy_kvcache = false; } @@ -185,7 +182,7 @@ void ov::npuw::LLMInferRequest::infer_generate(ov::SoPtr input_ids, m_kvcache_request->infer(); m_logits = m_kvcache_request->get_tensor(m_kvcache_out_ports.at("logits")); - m_kvcache_desc.num_stored_tokens += 1; + kvcache_desc.num_stored_tokens += 1; LOG_DEBUG("Write KV-cache for the new token to the correct input position for next iteration."); const std::size_t kStartOutputKVCacheLayers = 1u; @@ -194,13 +191,13 @@ 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) + const auto& kv_dim = (output_name.find("value") != std::string::npos && kvcache_desc.v_tensors_transposed) ? 3u - : m_kvcache_desc.dim; + : kvcache_desc.dim; auto kvcache_in_slice = make_tensor_slice(kvcache_in_tensor, kv_dim, - m_kvcache_desc.num_stored_tokens - 1, - m_kvcache_desc.num_stored_tokens); + kvcache_desc.num_stored_tokens - 1, + kvcache_desc.num_stored_tokens); auto kvcache_out_tensor = m_kvcache_request->get_tensor(m_kvcache_out_ports.at(output_name)); kvcache_out_tensor->copy_to(kvcache_in_slice._ptr); } diff --git a/src/plugins/intel_npu/src/plugin/npuw/llm_infer_request.hpp b/src/plugins/intel_npu/src/plugin/npuw/llm_infer_request.hpp index fbc6c702c4b62a..ea0b83d271b58e 100644 --- a/src/plugins/intel_npu/src/plugin/npuw/llm_infer_request.hpp +++ b/src/plugins/intel_npu/src/plugin/npuw/llm_infer_request.hpp @@ -15,8 +15,7 @@ namespace npuw { class LLMInferRequest final : public ov::ISyncInferRequest { public: - explicit LLMInferRequest(const std::shared_ptr& compiled_model, - const ov::npuw::LLMCompiledModel::KVCacheDesc& kvcache_desc); + explicit LLMInferRequest(const std::shared_ptr& compiled_model); void infer() override; @@ -44,7 +43,7 @@ class LLMInferRequest final : public ov::ISyncInferRequest { std::shared_ptr m_kvcache_request; std::shared_ptr m_prefill_request; - LLMCompiledModel::KVCacheDesc m_kvcache_desc; + std::shared_ptr m_npuw_llm_compiled_model; ov::SoPtr m_logits; bool m_need_copy_kvcache = false; diff --git a/src/plugins/intel_npu/src/utils/src/zero/CMakeLists.txt b/src/plugins/intel_npu/src/utils/src/zero/CMakeLists.txt index 76512189541525..5faf2c9f01c4a0 100644 --- a/src/plugins/intel_npu/src/utils/src/zero/CMakeLists.txt +++ b/src/plugins/intel_npu/src/utils/src/zero/CMakeLists.txt @@ -16,7 +16,7 @@ add_library(openvino::npu_zero_utils ALIAS ${TARGET_NAME}) set_target_properties(${TARGET_NAME} PROPERTIES EXPORT_NAME npu_zero_utils) add_library(level-zero-headers INTERFACE) -set_property(TARGET level-zero-headers APPEND PROPERTY INTERFACE_INCLUDE_DIRECTORIES $) +set_property(TARGET level-zero-headers APPEND PROPERTY INTERFACE_INCLUDE_DIRECTORIES $) add_library(LevelZero::Headers ALIAS level-zero-headers) target_include_directories(${TARGET_NAME} @@ -30,6 +30,7 @@ target_link_libraries(${TARGET_NAME} PUBLIC openvino::runtime::dev) # # targets install # + ov_install_static_lib(${TARGET_NAME} ${NPU_PLUGIN_COMPONENT}) ov_developer_package_export_targets(TARGET openvino::npu_zero_utils @@ -41,3 +42,14 @@ ov_install_static_lib(level-zero-ext ${NPU_PLUGIN_COMPONENT}) ov_developer_package_export_targets(TARGET level-zero-headers) ov_install_static_lib(level-zero-headers ${NPU_PLUGIN_COMPONENT}) + +if(TARGET ze_loader) + ov_developer_package_export_targets(TARGET ze_loader) + ov_install_static_lib(ze_loader ${NPU_PLUGIN_COMPONENT}) + + add_dependencies(${TARGET_NAME} ze_loader) + + # TODO: remove once https://github.com/oneapi-src/level-zero/pull/243 is merged and made a part of official release + ov_developer_package_export_targets(TARGET utils) + ov_install_static_lib(utils ${NPU_PLUGIN_COMPONENT}) +endif() diff --git a/tests/layer_tests/jax_tests/test_logistic.py b/tests/layer_tests/jax_tests/test_logistic.py new file mode 100644 index 00000000000000..19e5f0a81eba30 --- /dev/null +++ b/tests/layer_tests/jax_tests/test_logistic.py @@ -0,0 +1,37 @@ +# Copyright (C) 2018-2024 Intel Corporation +# SPDX-License-Identifier: Apache-2.0 + +import jax +import numpy as np +import pytest +from jax import numpy as jnp + +from jax_layer_test_class import JaxLayerTest + +rng = np.random.default_rng(5402) + + +class TestLogistic(JaxLayerTest): + def _prepare_input(self): + + input = jnp.array(np.random.uniform(-1000, 1000, self.input_shape).astype(self.input_type)) + return [input] + + def create_model(self, input_shape, input_type): + self.input_shape = input_shape + self.input_type = input_type + + def jax_logistic(input): + return jax.lax.logistic(input) + + return jax_logistic, None, 'logistic' + + @pytest.mark.parametrize("input_shape", [[2], [3, 4], [5,6,7]]) + @pytest.mark.parametrize("input_type", [np.float32, np.float64]) + @pytest.mark.nightly + @pytest.mark.precommit + @pytest.mark.precommit_jax_fe + def test_logistic(self, ie_device, precision, ir_version, input_shape, input_type): + self._test(*self.create_model(input_shape, input_type), + ie_device, precision, + ir_version) diff --git a/thirdparty/dependencies.cmake b/thirdparty/dependencies.cmake index 83f7e209bcf95b..51c9c63c971f05 100644 --- a/thirdparty/dependencies.cmake +++ b/thirdparty/dependencies.cmake @@ -80,7 +80,7 @@ if(ENABLE_INTEL_NPU) endif() endif() - if(NOT libze_loader_FOUND) + if(NOT level_zero_FOUND) add_subdirectory(thirdparty/level_zero EXCLUDE_FROM_ALL) add_library(LevelZero::LevelZero ALIAS ze_loader) endif()