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