natas_hw
/
openvino
0
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Merge branch 'master' into dependabot/pip/tests/pyyaml-6.0.1

This commit is contained in:
Ilya Lavrenov 2024-06-04 12:53:21 +04:00 committed by GitHub
parent b001a34e05 df6a25800d
commit 8b67a5d6f1
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
57 changed files with 2139 additions and 1038 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
docs/articles_en/documentation/openvino-ir-format/operation-sets/operation-specs/shape/squeeze-1.rst
View File

@ -17,7 +17,10 @@ Squeeze
**Detailed description**: *Squeeze* can be used with or without the second input tensor.
* If only the first input is provided, every dimension that is equal to 1 will be removed from it.
* With the second input provided, each value is an index of a dimension from the first tensor that is to be removed. Specified dimension has to be equal to 1, otherwise an error will be raised. Dimension indices can be specified directly, or by negative indices (counting dimensions from the end).
* With the second input provided, each value is an index of a dimension from the first tensor that is to be removed. Specified dimension should be equal to 1, otherwise it will be ignored and copied as is.
Dimension indices can be specified directly, or by negative indices (counting dimensions from the end).
Note: Updated behavior since 2024.3, request of squeezing dimension not equal to 1 is expected to be ignored instead of causing an error.
**Attributes**: *Squeeze* operation doesn't have attributes.
@ -87,4 +90,3 @@ Squeeze
</port>
</output>
</layer>

186
src/bindings/js/node/lib/addon.ts
View File

@ -21,46 +21,168 @@ type elementTypeString =
| 'f32'
| 'string';
/**
* Core represents an OpenVINO runtime Core entity.
*
* User applications can create several Core class instances,
* but in this case, the underlying plugins
* are created multiple times and not shared between several Core instances.
* It is recommended to have a single Core instance per application.
*/
interface Core {
/**
* Registers extensions to a Core object.
* @param libraryPath Path to the library with ov::Extension.
*/
addExtension(libraryPath: string): void;
/**
* Asynchronously creates a compiled model from a source {@link Model} object.
*
* You can create as many compiled models as needed and use them
* simultaneously (up to the limitation of the hardware resources).
* @param model The {@link Model} object acquired from {@link Core.readModel}
* @param deviceName The name of a device, to which the model is loaded.
* @param config An object with the key-value pairs
* (property name, property value): relevant only for this load operation.
*/
compileModel(
model: Model,
deviceName: string,
config?: { [option: string]: string }
config?: { [propertyName: string]: string }
): Promise<CompiledModel>;
/**
* Asynchronously reads a model and creates a compiled model
* from the IR/ONNX/PDPD file.
*
* This can be more efficient
* than using {@link Core.readModel} + core.compileModel(Model) flow
* especially for cases when caching is enabled and a cached model is
* available. You can create as many compiled models as needed and use
* them simultaneously (up to the limitation of the hardware resources).
* @param modelPath The path to a model.
* @param deviceName The name of a device, to which a model is loaded.
* @param config An object with the key-value pairs
* (property name, property value): relevant only for this load operation.
*/
compileModel(
modelPath: string,
deviceName: string,
config?: { [propertyName: string]: string }
): Promise<CompiledModel>;
/**
* A synchronous version of {@link Core.compileModel}.
* It creates a compiled model from a source model object.
*/
compileModelSync(
model: Model,
deviceName: string,
config?: { [option: string]: string }
config?: { [propertyName: string]: string }
): CompiledModel;
readModel(modelPath: string, weightsPath?: string): Promise<Model>;
readModel(
modelBuffer: Uint8Array, weightsBuffer?: Uint8Array): Promise<Model>;
readModelSync(modelPath: string, weightsPath?: string): Model;
readModelSync(modelBuffer: Uint8Array, weightsBuffer?: Uint8Array): Model;
importModelSync(modelStream: Buffer, device: string): CompiledModel;
importModelSync(
modelStream: Buffer,
device: string,
props: { [key: string]: string | number | boolean }
/**
* A synchronous version of {@link Core.compileModel}.
* It reads a model and creates a compiled model from the IR/ONNX/PDPD file.
*/
compileModelSync(
modelPath: string,
deviceName: string,
config?: { [propertyName: string]: string }
): CompiledModel;
/**
* It returns a list of available inference devices.
* Core objects go over all registered plugins.
* @returns The list of devices may include any of the following: CPU, GPU.0,
* GPU.1, NPU… If there is more than one device of a specific type, they are
* enumerated with .# suffix. Such enumerated devices can later be used
* as a device name in all Core methods, like compile_model, query_model,
* set_property and so on.
*/
getAvailableDevices(): string[];
/**
* It gets the properties dedicated to device behaviour.
* @param propertyName A property name.
*/
getProperty(propertyName: string): string | number | boolean;
/**
* It gets the properties dedicated to device behaviour.
* @param deviceName The name of a device, the properties of which you get.
* @param propertyName Property name.
*/
getProperty(
deviceName: string,
propertyName: string,
): string | number | boolean;
/**
* It returns information on the version of device plugins.
* @param deviceName A device name to identify a plugin.
*/
getVersions(deviceName: string): {
[deviceName: string]: {
buildNumber: string,
description: string,
},
};
setProperty(props: { [key: string]: string | number | boolean }): void;
/**
* It imports a previously exported compiled model.
* @param modelStream The input stream that contains a model, previously exported
* with the {@link CompiledModel.exportModelSync} method.
* @param device The name of a device, for which you import a compiled model.
* Note, if the device name was not used to compile the original model,
* an exception is thrown.
* @param config An object with the key-value pairs
* (property name, property value): relevant only for this load operation.
*/
importModelSync(
modelStream: Buffer,
device: string,
config?: { [key: string]: string | number | boolean }
): CompiledModel;
/**
* It reads models from the IR / ONNX / PDPD / TF and TFLite formats.
* @param modelPath The path to a model
* in the IR / ONNX / PDPD / TF or TFLite format.
* @param weightsPath The path to a data file for the IR format (.bin): if the path
* is empty, it tries to read the bin file with the same name as xml and if
* the bin file with the same name was not found, it loads IR without weights.
* For the ONNX format (.onnx), the weights parameter is not used.
* For the PDPD format (.pdmodel), the weights parameter is not used.
* For the TF format (.pb), the weights parameter is not used.
* For the TFLite format (*.tflite), the weights parameter is not used.
*/
readModel(modelPath: string, weightsPath?: string): Promise<Model>;
/**
* It reads models from the IR / ONNX / PDPD / TF and TFLite formats.
* @param modelBuffer Binary data with a model
* in the IR / ONNX / PDPD / TF or TFLite format.
* @param weightsBuffer Binary data with tensor data.
*/
readModel(
modelBuffer: Uint8Array, weightsBuffer?: Uint8Array): Promise<Model>;
/**
* A synchronous version of {@link Core.readModel}.
* It reads models from the IR / ONNX / PDPD / TF and TFLite formats.
*/
readModelSync(modelPath: string, weightsPath?: string): Model;
/**
* A synchronous version of {@link Core.readModel}.
* It reads models from the IR / ONNX / PDPD / TF and TFLite formats.
*/
readModelSync(modelBuffer: Uint8Array, weightsBuffer?: Uint8Array): Model;
/**
* It sets the properties.
* @param properties An object with the property name - property value pairs.
*/
setProperty(properties: { [key: string]: string | number | boolean }): void;
/**
* It sets the properties for a device.
* @param deviceName The name of a device.
* @param properties An object with the property name - property value pairs.
*/
setProperty(
deviceName: string,
props: { [key: string]: string | number | boolean },
properties: { [key: string]: string | number | boolean },
): void;
getProperty(propertyName: string): string | number | boolean,
getProperty(
deviceName: string,
propertyName: string,
): string | number | boolean;
addExtension(libraryPath: string): void;
}
interface CoreConstructor {
new(): Core;
@ -80,59 +202,59 @@ interface Model {
}
/**
* CompiledModel represents Model that is compiled for a specific device
* CompiledModel represents a model that is compiled for a specific device
* by applying multiple optimization transformations,
* then mapping to compute kernels.
*/
interface CompiledModel {
/** Gets all inputs of a compiled model. */
/** It gets all inputs of a compiled model. */
inputs: Output[];
/** Gets all outputs of a compiled model. */
/** It gets all outputs of a compiled model. */
outputs: Output[];
/**
* Creates an inference request object used to infer the compiled model.
* It creates an inference request object used to infer the compiled model.
* @return {InferRequest}
*/
createInferRequest(): InferRequest;
/**
* Exports the compiled model to binary data.
* It exports the compiled model to binary data.
* @remarks
* The exported model can be imported via the {@link Core.importModelSync}.
* @return {Buffer} The binary data that contains this compiled model.
* @return {Buffer} The binary data that contains the compiled model.
*/
exportModelSync(): Buffer;
/**
* Gets a single output of a compiled model.
* It gets a single output of a compiled model.
* If a model has more than one output, this method throws an exception.
* @returns {Output} A compiled model output.
*/
output(): Output;
/**
* Gets output of a compiled model identified by an index.
* It gets output of a compiled model identified by an index.
* @param index An output tensor index.
* @returns {Output} A compiled model output.
*/
output(index: number): Output;
/**
* Gets output of a compiled model identified by a tensorName.
* It gets output of a compiled model identified by a tensorName.
* @param name An output tensor name.
* @returns {Output} A compiled model output.
*/
output(name: string): Output;
/**
* Gets a single input of a compiled model.
* It gets a single input of a compiled model.
* If a model has more than one input, this method throws an exception.
* @returns {Output} A compiled model input.
*/
input(): Output;
/**
* Gets input of a compiled model identified by an index.
* It gets input of a compiled model identified by an index.
* @param index An input tensor index.
* @returns {Output} A compiled model input.
*/
input(index: number): Output;
/**
* Gets input of a compiled model identified by a tensorName.
* It gets input of a compiled model identified by a tensorName.
* @param name An input tensor name.
* @returns {Output} A compiled model input.
*/

12
src/common/transformations/include/transformations/utils/gen_pattern.hpp
View File

@ -1212,7 +1212,8 @@ public:
return false;
}
if (ele_type == ov::element::i32 || ele_type == ov::element::f32 || ele_type == ov::element::i64) {
if (ele_type == ov::element::i32 || ele_type == ov::element::i64 || ele_type == ov::element::f16 ||
ele_type == ov::element::f32) {
auto observed = constop->cast_vector<double>();
for (size_t i = 0; i < symbols.size(); i++)
detail::add_symbol_observed(sov, symbols[i], observed[i]);
@ -1259,6 +1260,15 @@ public:
}
}
if (pconst_node->get_output_element_type(0).is_real() &&
vconst_node->get_output_element_type(0).is_real()) {
auto p_values = pconst_node->cast_vector<float>();
auto v_values = vconst_node->cast_vector<float>();
if (p_values == v_values) {
continue;
}
}
_VERBOSE_LOG("expecting Constant of type ",
pconst_node->get_output_element_type(0),
" but got ",

15
src/common/transformations/src/transformations/common_optimizations/fuse_rotary_positional_embeddings.cpp
View File

@ -9,6 +9,7 @@
#include "itt.hpp"
#include "openvino/core/rt_info.hpp"
#include "openvino/op/util/shape_of_base.hpp"
#include "openvino/opsets/opset1.hpp"
#include "openvino/opsets/opset6.hpp"
#include "openvino/opsets/opset8.hpp"
@ -415,9 +416,9 @@ ov::pass::RoPEFusionGPTJ::RoPEFusionGPTJ() {
ov::pass::RoPEFusionChatGLM::RoPEFusionChatGLM(int split_output_id) {
MATCHER_SCOPE(RoPEFusionChatGLM);
auto qkv_linear = makePattern("f32[?,?,?]"); // f32[seq_length, batch_size, 4608]
auto qkv_linear = makePattern("[?,?,?]"); // [seq_length, batch_size, 4608]
auto seq_length = makePattern("i32[1]");
auto cos_sin_cache = makePattern("f32[?,?,?,?]"); // [max_pos_embeddings, batch_size, 32, 2]
auto cos_sin_cache = makePattern("[?,?,?,?]"); // [max_pos_embeddings, batch_size, 32, 2]
auto ndims = ov::gen_pattern::Symbol("ndims");
auto head_cnt = ov::gen_pattern::Symbol("head_cnt");
@ -538,9 +539,9 @@ ov::pass::RoPEFusionQwen::RoPEFusionQwen(int split_output_id) {
MATCHER_SCOPE(RoPEFusionQwen);
// rotary_emb_cos & rotary_emb_sin are sliced by present kv-length (past-kv-length + cur_len)
auto rotary_emb_cos = makePattern("f32[1,?,1,?]"); // [1,..4096,1,128]
auto rotary_emb_sin = makePattern("f32[1,?,1,?]"); // [1,..4096,1,128]
auto qkv_proj = makePattern("f32[?,?,?]"); // f32[?,?,12288]
auto rotary_emb_cos = makePattern("[1,?,1,?]"); // [1,..4096,1,128]
auto rotary_emb_sin = makePattern("[1,?,1,?]"); // [1,..4096,1,128]
auto qkv_proj = makePattern("[?,?,?]"); // [?,?,12288]
auto head_cnt = ov::gen_pattern::Symbol("head_cnt");
auto head_size = ov::gen_pattern::Symbol("head_size");
@ -559,8 +560,8 @@ ov::pass::RoPEFusionQwen::RoPEFusionQwen(int split_output_id) {
auto Multiply_567524 = makePattern<opset1::Multiply>({ShapeOf_485735, {-1}}, {{"auto_broadcast", "numpy"}});
auto Gather_377635 = makePattern<opset8::Gather>({Multiply_567524, {1}, 0}, {{"batch_dims", 0}});
auto input_ids = makePattern("i32[?,?]"); // [batch, length]
auto ShapeOf_409241 = makePattern<opset1::ShapeOf>({input_ids}, {});
auto input_ids = makePattern(); // [batch, length]
auto ShapeOf_409241 = makePattern<ov::op::util::ShapeOfBase>({input_ids}, {});
auto Gather_311651 = makePattern<opset8::Gather>({ShapeOf_409241, {1}, 0}, {{"batch_dims", 0}});
auto neg_Multiply = makePattern<opset1::Multiply>({Gather_311651, {-1}}, {{"auto_broadcast", "numpy"}});

6
src/core/shape_inference/include/scaled_dot_product_attention_shape_inference.hpp
View File

@ -73,9 +73,9 @@ std::vector<TRShape> shape_infer(const ScaledDotProductAttention* op,
const auto& attention_mask_rank = attention_mask.rank();
if (attention_mask_rank.is_static() && attention_mask_rank != 0) {
const auto& attention_mask_rank_len = attention_mask_rank.get_length();
bool attention_mask_input_correctness = attention_mask_rank_len >= 2 &&
DimType::merge(l_dim, l_dim, *(attention_mask.end() - 2)) &&
DimType::merge(s_dim, s_dim, *(attention_mask.end() - 1));
bool attention_mask_input_correctness =
attention_mask_rank_len >= 2 && DimType::broadcast_merge(l_dim, l_dim, *(attention_mask.end() - 2)) &&
DimType::broadcast_merge(s_dim, s_dim, *(attention_mask.end() - 1));
if (attention_mask_rank_len >= 3) {
attention_mask_input_correctness =
attention_mask_input_correctness &&

28
src/core/shape_inference/include/squeeze_shape_inference.hpp
View File

@ -54,17 +54,15 @@ std::vector<TRShape> shape_infer(const Squeeze* op,
unique_axes.reset(new std::set<int64_t>(axes->cbegin(), axes->cend()));
} else if (arg_rank.get_length() > 0 && shape_size(axes_shape.to_shape()) == 1) {
// The `axes` input is a single element tensor which is unique by definition, deducing output rank
NODE_VALIDATION_CHECK(op,
std::any_of(arg_shape.cbegin(),
arg_shape.cend(),
[](const DimType& dim) {
return dim.compatible(1);
}),
"Data input shape ",
arg_shape,
" doesn't contain squeezable dimension,"
" but axes input is expected to have one element.");
output_shape = PartialShape::dynamic(arg_rank.get_length() - 1);
const auto has_squeezable_dim =
std::any_of(arg_shape.cbegin(), arg_shape.cend(), [](const DimType& dim) {
return dim.compatible(1);
});
if (has_squeezable_dim) {
output_shape = PartialShape::dynamic(arg_rank.get_length() - 1);
} else {
output_shape = arg_shape;
}
return output_shapes;
}
}
@ -97,13 +95,11 @@ std::vector<TRShape> shape_infer(const Squeeze* op,
auto rm_axis_end = unique_axes->cend();
// Returns true if dimension not squeezable on axis from input axes.
const auto not_squeezable_at_axis = [&op, &rm_axis_iter, &rm_axis_end, &idx](const DimType& dim) {
const auto not_squeezable_at_axis = [&rm_axis_iter, &rm_axis_end, &idx](const DimType& dim) {
if ((rm_axis_iter != rm_axis_end) && (*rm_axis_iter == idx++)) {
NODE_VALIDATION_CHECK(op,
dim.compatible(1),
"provided axis value is invalid. Only axes of size 1 may be removed.");
++rm_axis_iter;
return false;
// Ignore: Pointed by axis, but not squeezable
return !dim.compatible(1);
} else {
return true;
}

12
src/core/tests/type_prop/squeeze.cpp
View File

@ -20,10 +20,10 @@ using namespace testing;
TEST(type_prop, squeeze_axes_invalid_value) {
auto param = make_shared<ov::op::v0::Parameter>(element::f32, Shape{1, 2, 3, 4});
auto axes_node = make_shared<ov::op::v0::Constant>(element::u64, Shape{2}, vector<int64_t>{0, 2});
const auto squeeze = std::make_shared<op::v0::Squeeze>(param, axes_node);
OV_EXPECT_THROW(auto s = make_shared<op::v0::Squeeze>(param, axes_node),
NodeValidationFailure,
HasSubstr("provided axis value is invalid. Only axes of size 1 may be removed."));
EXPECT_EQ(squeeze->get_element_type(), element::f32);
EXPECT_EQ(squeeze->get_output_partial_shape(0), (PartialShape{2, 3, 4}));
}
TEST(type_prop, squeeze_single_input) {
@ -53,10 +53,10 @@ TEST(type_prop, squeeze_incorrect_negative_axes) {
TEST(type_prop, squeeze_data_static_param_axes_1D_single_elem_static_shape_no_squeezable_dims) {
auto param = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, PartialShape{2, 2, 4});
const auto axes_node = std::make_shared<ov::op::v0::Parameter>(element::u64, PartialShape{1});
const auto squeeze = std::make_shared<op::v0::Squeeze>(param, axes_node);
OV_EXPECT_THROW(auto s = make_shared<op::v0::Squeeze>(param, axes_node),
NodeValidationFailure,
HasSubstr("doesn't contain squeezable dimension"));
EXPECT_EQ(squeeze->get_element_type(), element::f32);
EXPECT_EQ(squeeze->get_output_partial_shape(0), (PartialShape{2, 2, 4}));
}
TEST(type_prop, squeeze_data_static_param_axes_1D_two_elem_static_shape_squeezable_dims_two) {

14
src/plugins/intel_cpu/src/shape_inference/custom/reshape.cpp
View File

@ -83,30 +83,19 @@ Result SqueezeShapeInfer::infer(const std::vector<std::reference_wrapper<const V
ov::util::Cast<int64_t>());
std::vector<int64_t> originOutPattern = outPattern;
std::vector<bool> removeMask(inputShapeSize, false);
bool existError = false;
for (size_t i = 0; i < outputPatternSize; i++) {
if (outPattern[i] < 0) {
outPattern[i] = inputShapeSize + outPattern[i];
}
if (outPattern[i] >= 0 && outPattern[i] < static_cast<int64_t>(inputShapeSize)) {
if (outPattern[i] >= 0 && outPattern[i] < static_cast<int64_t>(inputShapeSize) && inputShape[outPattern[i]] == 1) {
removeMask[outPattern[i]] = true;
} else {
existError = true;
break;
}
}
for (size_t i = 0; i < inputShapeSize; i++) {
if (!removeMask[i]) {
outputShape.push_back(inputShape[i]);
} else if (inputShape[i] != 1) {
existError = true;
break;
}
}
if (existError) {
OPENVINO_THROW("[cpu]squeeze: the shape of input data ", ov::intel_cpu::vec2str(inputShape),
" conflicts with the squeeze pattern ", ov::intel_cpu::vec2str(originOutPattern));
}
} else {
for (size_t i = 0; i < inputShapeSize; i++) {
if (inputShape[i] != 1) {
@ -189,4 +178,3 @@ ShapeInferPtr ReshapeShapeInferFactory::makeShapeInfer() const {
} // namespace node
} // namespace intel_cpu
} // namespace ov

621
src/plugins/intel_cpu/tests/functional/custom/subgraph_tests/src/rotary_pos_emb.cpp
View File

@ -1,621 +0,0 @@
// Copyright (C) 2018-2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include <common_test_utils/ov_tensor_utils.hpp>
#include <openvino/opsets/opset1.hpp>
#include <openvino/opsets/opset8.hpp>
#include <string>
#include <tuple>
#include "common_test_utils/common_utils.hpp"
#include "shared_test_classes/base/ov_subgraph.hpp"
#include "utils/cpu_test_utils.hpp"
#include "utils/fusing_test_utils.hpp"
#include "transformations/utils/gen_pattern.hpp"
using namespace CPUTestUtils;
using namespace ov::gen_pattern;
using namespace ov;
namespace ov {
namespace test {
static ov::OutputVector makeCosSinCache(int max_position_embeddings, int rotary_ndims) {
std::vector<float> lut_sin(max_position_embeddings * rotary_ndims, 0.0f);
std::vector<float> lut_cos(max_position_embeddings * rotary_ndims, 0.0f);
// rotate_half style cos/sin table:
// y1 = cos(m*xita_i) * x1 - sin(m*xita_i) * x2
// y2 = cos(m*xita_i) * x2 + sin(m*xita_i) * x1
//
for (int i = 0, k = 0; i < rotary_ndims; i += 2, k++) {
auto xita_i = 1.0 / std::pow(10000.0, static_cast<double>(i) / rotary_ndims);
float* psin = lut_sin.data();
float* pcos = lut_cos.data();
for (int m = 0; m < max_position_embeddings; m++, psin += rotary_ndims, pcos += rotary_ndims) {
auto vsin = std::sin(xita_i * m);
auto vcos = std::cos(xita_i * m);
pcos[k] = pcos[k + rotary_ndims / 2] = vcos;
psin[k] = psin[k + rotary_ndims / 2] = vsin;
}
}
auto shape = ov::Shape({1, 1, static_cast<size_t>(max_position_embeddings), static_cast<size_t>(rotary_ndims)});
auto Cos = makeConst(ov::element::f32, shape, lut_cos);
auto Sin = makeConst(ov::element::f32, shape, lut_sin);
return {Cos, Sin};
}
static std::shared_ptr<ov::Model> buildROPE_Llama2(const int batch,
const int seq_length,
const int max_position_embeddings,
const int num_head,
const int ndims) {
auto input = std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{batch, -1, num_head, ndims});
auto pos_id_end = std::make_shared<ov::opset1::Parameter>(ov::element::i32, ov::Shape{});
auto pos_ids = std::make_shared<ov::opset1::Parameter>(ov::element::i32, PartialShape{1, -1});
auto cos_sin_cache = makeCosSinCache(max_position_embeddings, ndims);
auto Constant582 = cos_sin_cache[0];
auto Constant585 = cos_sin_cache[1];
// concat KV length
auto transpose_Transpose = makeOP<ov::op::v1::Transpose>({input, {0, 2, 1, 3}});
auto slice_Unsqueeze_426 = makeOP<ov::op::v0::Unsqueeze>({pos_id_end, 0});
auto ScatterUpdate_152236 = makeOP<ov::op::v3::ScatterUpdate>({{0, 0, 0}, {2}, slice_Unsqueeze_426, {0}});
auto slice_Slice = makeOP<ov::op::v1::StridedSlice>({Constant582, {0, 0, 0}, ScatterUpdate_152236, {1, 1, 1}},
{{"begin_mask", {1, 1, 0}},
{"end_mask", {1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto squeeze_Squeeze = makeOP<ov::op::v0::Squeeze>({slice_Slice, 1});
auto squeeze_Squeeze_435 = makeOP<ov::op::v0::Squeeze>({squeeze_Squeeze, 0});
auto index_441_Gather = makeOP<ov::op::v8::Gather>({squeeze_Squeeze_435, pos_ids, 0}, {{"batch_dims", 0}});
auto unsqueeze_Unsqueeze = makeOP<ov::op::v0::Unsqueeze>({index_441_Gather, 1});
auto mul_Multiply =
makeOP<ov::op::v1::Multiply>({transpose_Transpose, unsqueeze_Unsqueeze}, {{"auto_broadcast", "numpy"}});
auto size_ShapeOf_448 = makeOP<ov::op::v3::ShapeOf>({transpose_Transpose}, {{"output_type", "i32"}});
auto size_Gather_450 = makeOP<ov::op::v8::Gather>({size_ShapeOf_448, 3, 0}, {{"batch_dims", 0}});
auto floor_divide_Divide =
makeOP<ov::op::v1::Divide>({size_Gather_450, 2}, {{"auto_broadcast", "numpy"}, {"m_pythondiv", true}});
auto floor_divide_Floor = makeOP<ov::op::v0::Floor>({floor_divide_Divide});
auto slice_Unsqueeze_452 = makeOP<ov::op::v0::Unsqueeze>({floor_divide_Floor, 0});
auto ScatterUpdate_152312 = makeOP<ov::op::v3::ScatterUpdate>({{0, 0, 0, 0}, {3}, slice_Unsqueeze_452, {0}});
auto slice_Slice_459 = makeOP<ov::op::v1::StridedSlice>(
{transpose_Transpose, ScatterUpdate_152312, {0ll, 0ll, 0ll, LLONG_MAX}, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto Constant_182988 = makeConst(element::f32,
ov::Shape({
1,
1,
1,
1,
}),
{-1.000000f});
auto neg_Multiply = makeOP<ov::op::v1::Multiply>({slice_Slice_459, Constant_182988}, {{"auto_broadcast", "numpy"}});
auto ScatterUpdate_152368 = makeOP<ov::op::v3::ScatterUpdate>({{0, 0, 0, 0}, {3}, slice_Unsqueeze_452, {0}});
auto slice_Slice2 =
makeOP<ov::op::v1::StridedSlice>({transpose_Transpose, {0, 0, 0, 0}, ScatterUpdate_152368, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto cat_Concat = makeOP<ov::op::v0::Concat>({neg_Multiply, slice_Slice2}, {{"axis", -1}});
auto ScatterUpdate_152421 = makeOP<ov::op::v3::ScatterUpdate>({{0, 0, 0}, {2}, slice_Unsqueeze_426, {0}});
auto slice_Slice_433 = makeOP<ov::op::v1::StridedSlice>({Constant585, {0, 0, 0}, ScatterUpdate_152421, {1, 1, 1}},
{{"begin_mask", {1, 1, 0}},
{"end_mask", {1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto squeeze_Squeeze_436 = makeOP<ov::op::v0::Squeeze>({slice_Slice_433, 1});
auto squeeze_Squeeze_437 = makeOP<ov::op::v0::Squeeze>({squeeze_Squeeze_436, 0});
auto index_446_Gather = makeOP<ov::op::v8::Gather>({squeeze_Squeeze_437, pos_ids, 0}, {{"batch_dims", 0}});
auto unsqueeze_Unsqueeze_447 = makeOP<ov::op::v0::Unsqueeze>({index_446_Gather, 1});
auto mul_Multiply_463 =
makeOP<ov::op::v1::Multiply>({cat_Concat, unsqueeze_Unsqueeze_447}, {{"auto_broadcast", "numpy"}});
auto add_Add = makeOP<ov::op::v1::Add>({mul_Multiply, mul_Multiply_463}, {{"auto_broadcast", "numpy"}});
return std::make_shared<ov::Model>(ov::NodeVector{add_Add}, ov::ParameterVector{input, pos_id_end, pos_ids});
}
class RoPECPUTestLlama2 : public SubgraphBaseTest {
public:
ov::Tensor create_i32_tensor(const ov::Shape& shape, int start, int step = 1) {
auto tensor = ov::Tensor(ov::element::i32, shape);
auto* ptr = static_cast<int32_t*>(tensor.data());
for (size_t i = 0; i < tensor.get_size(); i++) {
ptr[i] = start;
start += step;
}
return tensor;
}
void generate_inputs(const std::vector<ov::Shape>& targetInputStaticShapes) override {
const auto& funcInputs = function->inputs();
const int position_id_start = 15;
auto& input_shape = targetInputStaticShapes[0];
auto seq_length = input_shape[1];
ov::test::utils::InputGenerateData in_data;
in_data.start_from = -1;
in_data.range = 2;
in_data.resolution = 32768;
ov::Tensor t_input = utils::create_and_fill_tensor(funcInputs[0].get_element_type(), input_shape, in_data);
ov::Tensor t_position_id_end = create_i32_tensor(ov::Shape({}), position_id_start + seq_length);
ov::Tensor t_position_ids = create_i32_tensor(ov::Shape({1, seq_length}), position_id_start);
inputs.clear();
inputs.insert({funcInputs[0].get_node_shared_ptr(), t_input});
inputs.insert({funcInputs[1].get_node_shared_ptr(), t_position_id_end});
inputs.insert({funcInputs[2].get_node_shared_ptr(), t_position_ids});
}
protected:
void SetUp() override {
targetDevice = ov::test::utils::DEVICE_CPU;
const int batch = 2;
const int seq_length = 7;
const size_t max_position_embeddings = 2048;
const size_t ndims = 128;
const size_t num_head = 32;
InputShape inpShape = {{batch, seq_length, num_head, ndims}, {{batch, seq_length, num_head, ndims}}};
init_input_shapes({inpShape});
function = buildROPE_Llama2(batch, seq_length, max_position_embeddings, num_head, ndims);
}
};
TEST_F(RoPECPUTestLlama2, smoke_CompareWithRefs) {
run();
CheckNumberOfNodesWithType(compiledModel, "RoPE", 1);
}
class RoPECPUTestChatGLM : public SubgraphBaseTest {
public:
ov::Tensor create_i32_tensor(const ov::Shape& shape, int start, int step = 1) {
auto tensor = ov::Tensor(ov::element::i32, shape);
auto* ptr = static_cast<int32_t*>(tensor.data());
for (size_t i = 0; i < tensor.get_size(); i++) {
ptr[i] = start;
start += step;
}
return tensor;
}
void generate_inputs(const std::vector<ov::Shape>& targetInputStaticShapes) override {
const auto& funcInputs = function->inputs();
auto& input_shape = targetInputStaticShapes[0];
auto seq_length = input_shape[0];
// auto batch = input_shape[1];
ov::Tensor t_input =
utils::create_and_fill_tensor(funcInputs[0].get_element_type(), input_shape, 2, -1.0f, 32768);
ov::Tensor t_cos_sin_cache =
utils::create_and_fill_tensor(funcInputs[1].get_element_type(), {32768, 32, 2}, 2, -1.0f, 32768);
ov::Tensor t_position_ids = create_i32_tensor(ov::Shape({1, seq_length}), 15);
inputs.clear();
inputs.insert({funcInputs[0].get_node_shared_ptr(), t_input});
inputs.insert({funcInputs[1].get_node_shared_ptr(), t_cos_sin_cache});
inputs.insert({funcInputs[2].get_node_shared_ptr(), t_position_ids});
}
protected:
std::shared_ptr<ov::Model> buildROPE_ChatGLM(int batch, int head_cnt, int rotary_dims) {
auto input =
std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{-1, batch, 4096 + 256 + 256});
auto cos_sin_cache = std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{32768, 32, 2});
auto position_ids = std::make_shared<ov::opset1::Parameter>(ov::element::i32, PartialShape{-1, -1});
auto __module_transformer_index_67_Gather =
makeOP<opset8::Gather>({cos_sin_cache, position_ids, 0}, {{"batch_dims", 0}});
auto __module_transformer_transpose_Transpose =
makeOP<opset1::Transpose>({__module_transformer_index_67_Gather, {1, 0, 2, 3}});
auto size_ShapeOf_110 =
makeOP<opset3::ShapeOf>({__module_transformer_transpose_Transpose}, {{"output_type", "i32"}});
auto __getitem___Gather = makeOP<opset8::Gather>({size_ShapeOf_110, -2, 0}, {{"batch_dims", 0}});
auto mul_Multiply = makeOP<opset1::Multiply>({__getitem___Gather, 2}, {{"auto_broadcast", "numpy"}});
auto slice_Unsqueeze_112 = makeOP<opset1::Unsqueeze>({mul_Multiply, 0});
auto floordiv_Divide =
makeOP<opset1::Divide>({mul_Multiply, 2}, {{"auto_broadcast", "numpy"}, {"m_pythondiv", true}});
auto floordiv_Floor = makeOP<opset1::Floor>({floordiv_Divide});
auto ListConstruct_126_Reshape_2 = makeOP<opset1::Reshape>({floordiv_Floor, {-1}}, {{"special_zero", false}});
auto ListUnpack_321 = makeOP<opset1::VariadicSplit>({input, -1, {4096, 256, 256}});
auto view_Reshape =
makeOP<opset1::Reshape>({ListUnpack_321->output(0), {0, 0, 32, 128}}, {{"special_zero", true}});
auto ScatterUpdate_229053 = makeOP<opset3::ScatterUpdate>({{0, 0, 0, 0}, {3}, slice_Unsqueeze_112, {0}});
auto slice_Slice_357 =
makeOP<opset1::StridedSlice>({view_Reshape, {0, 0, 0, 0}, ScatterUpdate_229053, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto size_ShapeOf_346 = makeOP<opset3::ShapeOf>({view_Reshape}, {{"output_type", "i32"}});
auto size_Gather_348 = makeOP<opset8::Gather>({size_ShapeOf_346, 0, 0}, {{"batch_dims", 0}});
auto ListConstruct_372_Reshape = makeOP<opset1::Reshape>({size_Gather_348, {-1}}, {{"special_zero", false}});
auto size_Gather_351 = makeOP<opset8::Gather>({size_ShapeOf_346, {2}, 0}, {{"batch_dims", 0}});
auto ListConstruct_372_Concat =
makeOP<opset1::Concat>({ListConstruct_372_Reshape, {-1}, size_Gather_351, ListConstruct_126_Reshape_2, {2}},
{{"axis", 0}});
auto reshape_Reshape_373 =
makeOP<opset1::Reshape>({slice_Slice_357, ListConstruct_372_Concat}, {{"special_zero", false}});
auto select_Gather_381 = makeOP<opset8::Gather>({reshape_Reshape_373, 0, -1}, {{"batch_dims", 0}});
auto slice_Unsqueeze_367 = makeOP<opset1::Unsqueeze>({size_Gather_348, 0});
auto slice_Slice_369 =
makeOP<opset1::StridedSlice>({__module_transformer_transpose_Transpose, {0}, slice_Unsqueeze_367, {1}},
{{"begin_mask", {0}},
{"end_mask", {0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto size_ShapeOf_374 = makeOP<opset3::ShapeOf>({reshape_Reshape_373}, {{"output_type", "i32"}});
auto size_Gather_376 = makeOP<opset8::Gather>({size_ShapeOf_374, {3}, 0}, {{"batch_dims", 0}});
auto ListConstruct_379_Concat =
makeOP<opset1::Concat>({ListConstruct_372_Reshape, {-1}, {1}, size_Gather_376, {2}}, {{"axis", 0}});
auto view_Reshape_380 =
makeOP<opset1::Reshape>({slice_Slice_369, ListConstruct_379_Concat}, {{"special_zero", false}});
auto select_Gather_382 = makeOP<opset8::Gather>({view_Reshape_380, 0, -1}, {{"batch_dims", 0}});
auto mul_Multiply_383 =
makeOP<opset1::Multiply>({select_Gather_381, select_Gather_382}, {{"auto_broadcast", "numpy"}});
auto select_Gather_384 = makeOP<opset8::Gather>({reshape_Reshape_373, 1, -1}, {{"batch_dims", 0}});
auto select_Gather_385 = makeOP<opset8::Gather>({view_Reshape_380, 1, -1}, {{"batch_dims", 0}});
auto mul_Multiply_386 =
makeOP<opset1::Multiply>({select_Gather_384, select_Gather_385}, {{"auto_broadcast", "numpy"}});
auto sub_Subtract_389 =
makeOP<opset1::Subtract>({mul_Multiply_383, mul_Multiply_386}, {{"auto_broadcast", "numpy"}});
auto Unsqueeze_62716 = makeOP<opset1::Unsqueeze>({sub_Subtract_389, -1});
auto mul_Multiply_391 =
makeOP<opset1::Multiply>({select_Gather_384, select_Gather_382}, {{"auto_broadcast", "numpy"}});
auto mul_Multiply_393 =
makeOP<opset1::Multiply>({select_Gather_381, select_Gather_385}, {{"auto_broadcast", "numpy"}});
auto add_Add_396 = makeOP<opset1::Add>({mul_Multiply_391, mul_Multiply_393}, {{"auto_broadcast", "numpy"}});
auto Unsqueeze_62717 = makeOP<opset1::Unsqueeze>({add_Add_396, -1});
auto stack_401 = makeOP<opset1::Concat>({Unsqueeze_62716, Unsqueeze_62717}, {{"axis", -1}});
auto flatten_ShapeOf_402 = makeOP<opset3::ShapeOf>({stack_401}, {{"output_type", "i32"}});
auto flatten_Slice_417 = makeOP<opset1::StridedSlice>({flatten_ShapeOf_402, {0}, {3}, {1}},
{{"begin_mask", {0}},
{"end_mask", {0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto flatten_Concat_420 = makeOP<opset1::Concat>({flatten_Slice_417, {-1}}, {{"axis", 0}});
auto flatten_Reshape_421 = makeOP<opset1::Reshape>({stack_401, flatten_Concat_420}, {{"special_zero", true}});
auto ScatterUpdate_229067 = makeOP<opset3::ScatterUpdate>({{0, 0, 0, 0}, {3}, slice_Unsqueeze_112, {0}});
auto slice_Slice_363 =
makeOP<opset1::StridedSlice>({view_Reshape, ScatterUpdate_229067, {0, 0, 0, INT_MAX}, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto cat_Concat_425 = makeOP<opset1::Concat>({flatten_Reshape_421, slice_Slice_363}, {{"axis", -1}});
return std::make_shared<ov::Model>(ov::NodeVector{cat_Concat_425},
ov::ParameterVector{input, cos_sin_cache, position_ids});
}
void SetUp() override {
targetDevice = ov::test::utils::DEVICE_CPU;
const int batch = 2;
const int seq_length = 7;
const int num_head = 32;
const int rotary_dims = 64;
InputShape inpShape = {{-1, batch, 4096 + 256 + 256}, {{seq_length, batch, 4096 + 256 + 256}}};
init_input_shapes({inpShape});
function = buildROPE_ChatGLM(batch, num_head, rotary_dims);
}
};
TEST_F(RoPECPUTestChatGLM, smoke_CompareWithRefs) {
run();
CheckNumberOfNodesWithType(compiledModel, "RoPE", 1);
}
class RoPECPUTestQwen7b : public SubgraphBaseTest, public testing::WithParamInterface<bool> {
public:
static std::string getTestCaseName(const testing::TestParamInfo<bool>& obj) {
const bool specialReshape = obj.param;
std::ostringstream result;
result << "specialReshape=" << specialReshape << std::endl;
return result.str();
}
void generate_inputs(const std::vector<ov::Shape>& targetInputStaticShapes) override {
const auto& funcInputs = function->inputs();
auto& input_shape = targetInputStaticShapes[0];
ov::Tensor t_input =
utils::create_and_fill_tensor(funcInputs[0].get_element_type(), input_shape, 2, -1.0f, 32768);
ov::Tensor t_cos_cache =
utils::create_and_fill_tensor(funcInputs[1].get_element_type(), {1, 4096, 1, 128}, 2, -1.0f, 32768);
ov::Tensor t_sin_cache =
utils::create_and_fill_tensor(funcInputs[1].get_element_type(), {1, 4096, 1, 128}, 2, -1.0f, 32768);
inputs.clear();
inputs.insert({funcInputs[0].get_node_shared_ptr(), t_input});
inputs.insert({funcInputs[1].get_node_shared_ptr(), t_cos_cache});
inputs.insert({funcInputs[2].get_node_shared_ptr(), t_sin_cache});
}
protected:
std::shared_ptr<ov::Model> buildROPE_QWen7b(bool specialReshape) {
auto input =
std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{-1, -1, 4096 + 4096 + 4096});
auto cos_cache = std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{1, -1, 1, 128});
auto sin_cache = std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{1, -1, 1, 128});
auto ListUnpack_389_VariadicSplit = makeOP<opset1::VariadicSplit>({input, 2, {4096, 4096, -1}});
auto view_Reshape = makeOP<opset1::Reshape>({ListUnpack_389_VariadicSplit->output(0), {0, 0, 32, 128}},
{{"special_zero", true}});
auto size_ShapeOf_414 = makeOP<opset3::ShapeOf>({view_Reshape}, {{"output_type", "i32"}});
auto size_Gather_416 = makeOP<opset8::Gather>({size_ShapeOf_414, 1, 0}, {{"batch_dims", 0}});
auto neg_Multiply = makeOP<opset1::Multiply>({size_Gather_416, -1}, {{"auto_broadcast", "numpy"}});
auto slice_Unsqueeze_422 = makeOP<opset1::Unsqueeze>({neg_Multiply, 0});
auto ScatterUpdate_261437 = makeOP<opset3::ScatterUpdate>({{0, 0}, {1}, slice_Unsqueeze_422, {0}});
auto slice_Slice_425 = makeOP<opset1::StridedSlice>({cos_cache, ScatterUpdate_261437, {0ll, LLONG_MAX}, {1, 1}},
{{"begin_mask", {1, 0}},
{"end_mask", {1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto slice_Slice_431 =
makeOP<opset1::StridedSlice>({slice_Slice_425, {0, 0, 0}, {0ll, 0ll, LLONG_MAX}, {1, 1, 1}},
{{"begin_mask", {1, 1, 0}},
{"end_mask", {1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto slice_Slice_437 =
makeOP<opset1::StridedSlice>({slice_Slice_431, {0, 0, 0, 0}, {0ll, 0ll, 0ll, LLONG_MAX}, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto size_ShapeOf_462 = makeOP<opset3::ShapeOf>({slice_Slice_437}, {{"output_type", "i32"}});
auto size_Gather_464 = makeOP<opset8::Gather>({size_ShapeOf_462, {3}, 0}, {{"batch_dims", 0}});
auto ScatterUpdate_261533 = makeOP<opset3::ScatterUpdate>({{0, 0, 0, 0}, {3}, size_Gather_464, {0}});
auto slice_Slice_470 =
makeOP<opset1::StridedSlice>({view_Reshape, {0, 0, 0, 0}, ScatterUpdate_261533, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto mul_Multiply = makeOP<opset1::Multiply>({slice_Slice_470, slice_Slice_437}, {{"auto_broadcast", "numpy"}});
auto size_ShapeOf_478 = makeOP<opset3::ShapeOf>({slice_Slice_470}, {{"output_type", "i32"}});
auto Gather_239390 = makeOP<opset8::Gather>({size_ShapeOf_478, {0, 1, 2}, 0}, {{"batch_dims", 0}});
auto size_Gather_489 = makeOP<opset8::Gather>({size_ShapeOf_478, 3, 0}, {{"batch_dims", 0}});
auto floor_divide_Divide =
makeOP<opset1::Divide>({size_Gather_489, 2}, {{"auto_broadcast", "numpy"}, {"m_pythondiv", true}});
auto floor_divide_Floor = makeOP<opset1::Floor>({floor_divide_Divide});
auto ListConstruct_493_Reshape_3 =
makeOP<opset1::Reshape>({floor_divide_Floor, {-1}}, {{"special_zero", false}});
auto ListConstruct_493_Concat =
makeOP<opset1::Concat>({Gather_239390, {2}, ListConstruct_493_Reshape_3}, {{"axis", 0}});
std::shared_ptr<ov::Node> reshape_Reshape = nullptr;
if (specialReshape) {
reshape_Reshape = makeOP<opset1::Reshape>({slice_Slice_470, {0, 0, 32, 2, 64}}, {{"special_zero", true}});
} else {
reshape_Reshape =
makeOP<opset1::Reshape>({slice_Slice_470, ListConstruct_493_Concat}, {{"special_zero", false}});
}
auto ListUnpack_496_Split = makeOP<opset1::Split>({reshape_Reshape, -2}, {{"num_splits", 2}});
auto ListUnpack_496_Squeeze_0 = makeOP<opset1::Squeeze>({ListUnpack_496_Split->output(1), -2});
auto Constant_296840_compressed = makeConst(element::f16,
ov::Shape({
1,
1,
1,
1,
}),
{-1});
auto Constant_296840 = makeOP<opset1::Convert>({Constant_296840_compressed}, {{"destination_type", "f32"}});
auto neg_Multiply_499 =
makeOP<opset1::Multiply>({ListUnpack_496_Squeeze_0, Constant_296840}, {{"auto_broadcast", "numpy"}});
auto ListUnpack_496_Squeeze = makeOP<opset1::Squeeze>({ListUnpack_496_Split->output(0), -2});
auto cat_Concat = makeOP<opset1::Concat>({neg_Multiply_499, ListUnpack_496_Squeeze}, {{"axis", -1}});
auto slice_Slice_449 = makeOP<opset1::StridedSlice>({sin_cache, ScatterUpdate_261437, {0ll, LLONG_MAX}, {1, 1}},
{{"begin_mask", {1, 0}},
{"end_mask", {1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto slice_Slice_455 =
makeOP<opset1::StridedSlice>({slice_Slice_449, {0, 0, 0}, {0ll, 0ll, LLONG_MAX}, {1, 1, 1}},
{{"begin_mask", {1, 1, 0}},
{"end_mask", {1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto slice_Slice_461 =
makeOP<opset1::StridedSlice>({slice_Slice_455, {0, 0, 0, 0}, {0ll, 0ll, 0ll, LLONG_MAX}, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto mul_Multiply_503 = makeOP<opset1::Multiply>({cat_Concat, slice_Slice_461}, {{"auto_broadcast", "numpy"}});
auto add_Add = makeOP<opset1::Add>({mul_Multiply, mul_Multiply_503}, {{"auto_broadcast", "numpy"}});
return std::make_shared<ov::Model>(ov::NodeVector{add_Add}, ov::ParameterVector{input, cos_cache, sin_cache});
}
void SetUp() override {
targetDevice = ov::test::utils::DEVICE_CPU;
const bool specialReshape = this->GetParam();
const int batch = 2;
const int seq_length = 7;
InputShape inpShape = {{batch, -1, 4096 + 4096 + 4096}, {{batch, seq_length, 4096 + 4096 + 4096}}};
init_input_shapes({inpShape});
function = buildROPE_QWen7b(specialReshape);
}
};
TEST_P(RoPECPUTestQwen7b, smoke_CompareWithRefs) {
run();
CheckNumberOfNodesWithType(compiledModel, "RoPE", 1);
}
INSTANTIATE_TEST_SUITE_P(smoke_RoPECPUTestQwen7b,
RoPECPUTestQwen7b,
::testing::Values(true, false),
RoPECPUTestQwen7b::getTestCaseName);
class RoPECPUTestGPTJ : public SubgraphBaseTest, public testing::WithParamInterface<bool> {
public:
static std::string getTestCaseName(const testing::TestParamInfo<bool>& obj) {
bool hasShapeOf;
hasShapeOf = obj.param;
std::ostringstream result;
result << "hasShapeOf=" << hasShapeOf << std::endl;
return result.str();
}
void generate_inputs(const std::vector<ov::Shape>& targetInputStaticShapes) override {
const auto& funcInputs = function->inputs();
auto& input_shape = targetInputStaticShapes[0];
auto& sincos_shape = targetInputStaticShapes[1];
ov::Tensor t_input =
utils::create_and_fill_tensor(funcInputs[0].get_element_type(), input_shape, 2, -1.0f, 32768);
ov::Tensor t_cos_sin_cache =
utils::create_and_fill_tensor(funcInputs[1].get_element_type(), sincos_shape, 2, -1.0f, 32768);
inputs.clear();
inputs.insert({funcInputs[0].get_node_shared_ptr(), t_input});
inputs.insert({funcInputs[1].get_node_shared_ptr(), t_cos_sin_cache});
}
protected:
std::shared_ptr<ov::Model> buildROPE_GPTJ(const int num_head,
const int hidden_dims,
const int rotary_dims,
bool hasShapeOf) {
auto int32_max = std::numeric_limits<std::int32_t>::max();
auto input =
std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{-1, -1, num_head, hidden_dims});
auto sincos = std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{-1, -1, rotary_dims});
auto slice_Slice_965 =
makeOP<ov::op::v1::StridedSlice>({input, {0, 0, 0, 0}, {0, 0, 0, rotary_dims}, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
slice_Slice_965->set_friendly_name("slice_Slice_965");
auto varsplit = makeOP<ov::op::v1::VariadicSplit>({sincos, -1, {rotary_dims / 2, -1}});
varsplit->set_output_size(2);
varsplit->set_friendly_name("varsplit");
auto unsqueeze_sin = makeOP<opset1::Unsqueeze>({varsplit->output(0), 2});
auto unsqueeze_cos = makeOP<opset1::Unsqueeze>({varsplit->output(1), 2});
std::vector<int32_t> gather_idx(rotary_dims, 1);
int32_t v = 0;
for (size_t i = 0; i < gather_idx.size(); i += 2, v++) {
gather_idx[i] = v;
gather_idx[i + 1] = v;
}
auto const_idx = makeConst(ov::element::i32, ov::Shape({static_cast<size_t>(rotary_dims)}), gather_idx);
auto constant_155588 = makeConst(element::f32,
ov::Shape({
1,
1,
1,
1,
}),
{-1.000000f});
auto repeat_interleave_sin = makeOP<opset8::Gather>({unsqueeze_sin, const_idx, 3}, {{"batch_dims", 0}});
auto repeat_interleave_cos = makeOP<opset8::Gather>({unsqueeze_cos, const_idx, 3}, {{"batch_dims", 0}});
repeat_interleave_sin->set_friendly_name("repeat_interleave_sin");
repeat_interleave_cos->set_friendly_name("repeat_interleave_cos");
// x interleave (-x[:,:,:, 1::2], x[:,:,:, 0::2])
auto slice_Slice_1174 =
makeOP<ov::op::v1::StridedSlice>({slice_Slice_965, {0, 0, 0, 1}, {0, 0, 0, int32_max}, {1, 1, 1, 2}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto neg_Multiply_1177 =
makeOP<opset1::Multiply>({slice_Slice_1174, constant_155588}, {{"auto_broadcast", "numpy"}});
auto Unsqueeze_65524 = makeOP<opset1::Unsqueeze>({neg_Multiply_1177, -1});
auto slice_Slice_1168 =
makeOP<ov::op::v1::StridedSlice>({slice_Slice_965, {0, 0, 0, 0}, {0, 0, 0, int32_max}, {1, 1, 1, 2}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto Unsqueeze_65525 = makeOP<opset1::Unsqueeze>({slice_Slice_1168, -1});
auto stack_1182 = makeOP<opset1::Concat>({Unsqueeze_65524, Unsqueeze_65525}, {{"axis", -1}});
auto flatten_Reshape_1198 =
makeOP<opset1::Reshape>({stack_1182, {0, 0, num_head, rotary_dims}}, {{"special_zero", true}});
// x*cos [B,L,H,ndims]
auto mul_cos =
makeOP<opset1::Multiply>({slice_Slice_965, repeat_interleave_cos}, {{"auto_broadcast", "numpy"}});
mul_cos->set_friendly_name("mul_cos");
auto mul_sin =
makeOP<opset1::Multiply>({flatten_Reshape_1198, repeat_interleave_sin}, {{"auto_broadcast", "numpy"}});
// *cos + *sin
auto rotary_emb = makeOP<opset1::Add>({mul_cos, mul_sin}, {{"auto_broadcast", "numpy"}});
auto slice_Slice_971 =
makeOP<ov::op::v1::StridedSlice>({input, {0, 0, 0, rotary_dims}, {0, 0, 0, int32_max}, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto cat_Concat_1211 = makeOP<opset1::Concat>({rotary_emb, slice_Slice_971}, {{"axis", -1}});
auto permute_Transpose_1213 = makeOP<opset1::Transpose>({cat_Concat_1211, {0, 2, 1, 3}});
ov::NodeVector model_output = {permute_Transpose_1213};
if (hasShapeOf) {
auto shapeOf = makeOP<opset1::ShapeOf>({rotary_emb}, {{"output_type", "i32"}});
auto gather = makeOP<opset8::Gather>({shapeOf, {1}, 0}, {{"batch_dims", 0}});
model_output.push_back(gather);
}
return std::make_shared<ov::Model>(model_output, ov::ParameterVector{input, sincos});
}
void SetUp() override {
targetDevice = ov::test::utils::DEVICE_CPU;
bool hasShapeOf = this->GetParam();
const int batch = 2;
const int seq_length = 7;
const int num_head = 16;
const int hidden_dims = 256;
const int rotary_dims = 64;
InputShape input = {{batch, seq_length, num_head, hidden_dims}, {{batch, seq_length, num_head, hidden_dims}}};
InputShape sincos = {{batch, seq_length, rotary_dims}, {{batch, seq_length, rotary_dims}}};
init_input_shapes({input, sincos});
function = buildROPE_GPTJ(num_head, hidden_dims, rotary_dims, hasShapeOf);
}
};
TEST_P(RoPECPUTestGPTJ, smoke_CompareWithRefs) {
run();
CheckNumberOfNodesWithType(compiledModel, "RoPE", 1);
}
INSTANTIATE_TEST_SUITE_P(smoke_RoPECPUTestGPTJ,
RoPECPUTestGPTJ,
::testing::Values(true, false),
RoPECPUTestGPTJ::getTestCaseName);
} // namespace test
} // namespace ov

2
src/plugins/intel_cpu/tests/functional/shared_tests_instances/skip_tests_config.cpp
View File

@ -371,6 +371,8 @@ std::vector<std::string> disabledTestPatterns() {
retVector.emplace_back(R"(smoke_VariableState/OVInferRequestVariableStateTest.*)");
// Issue: 141705
retVector.emplace_back(R"(.*smoke_arm_Deconv_2D_Planar_FP16/DeconvolutionLayerCPUTest.*INFERENCE_PRECISION_HINT=f16.*)");
retVector.emplace_back(R"(.*smoke_RoPETest.*)");
#endif
#if defined(OPENVINO_ARCH_ARM)

32
src/plugins/intel_cpu/tests/functional/shared_tests_instances/subgraph_tests/rotary_pos_emb.cpp Normal file
View File

@ -0,0 +1,32 @@
// Copyright (C) 2018-2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "subgraph_tests/rotary_pos_emb.hpp"
namespace ov {
namespace test {
INSTANTIATE_TEST_SUITE_P(smoke_RoPETestLlama2,
RoPETestLlama2,
::testing::Values(ov::test::utils::DEVICE_CPU),
RoPETestLlama2::getTestCaseName);
INSTANTIATE_TEST_SUITE_P(smoke_RoPETestChatGLM,
RoPETestChatGLM,
::testing::Values(ov::test::utils::DEVICE_CPU),
RoPETestChatGLM::getTestCaseName);
INSTANTIATE_TEST_SUITE_P(smoke_RoPETestQwen7b,
RoPETestQwen7b,
::testing::Combine(::testing::Values(true, false),
::testing::Values(ov::test::utils::DEVICE_CPU)),
RoPETestQwen7b::getTestCaseName);
INSTANTIATE_TEST_SUITE_P(smoke_RoPETestGPTJ,
RoPETestGPTJ,
::testing::Combine(::testing::Values(true, false),
::testing::Values(ov::test::utils::DEVICE_CPU)),
RoPETestGPTJ::getTestCaseName);
} // namespace test
} // namespace ov

58
src/plugins/intel_cpu/tests/unit/shape_inference_test/custom_shape_infer/squeeze.cpp
View File

@ -87,50 +87,30 @@ INSTANTIATE_TEST_SUITE_P(
make_tuple(unit_test::ShapeVector{{2, 6, 7, 8, 1}, {2}}, std::vector<int64_t>{-1, -1}, StaticShape({2, 6, 7, 8}))),
SqueezeCpuShapeInferenceTest::getTestCaseName);
using SqueezeCpuShapeInferenceThrowExceptionTest = SqueezeCpuShapeInferenceTest;
TEST_P(SqueezeCpuShapeInferenceThrowExceptionTest, wrong_pattern) {
// Tests with non-squeezable dims pointed by axes (no throw, ignore)
class SqueezeCpuShapeInferenceTestNonSqueeezable : public SqueezeCpuShapeInferenceTest {};
TEST_P(SqueezeCpuShapeInferenceTestNonSqueeezable, shape_inference_non_squeezable_with_const_map) {
const auto axes_node = std::make_shared<op::v0::Parameter>(element::i64, PartialShape::dynamic());
const auto op = make_op(arg, axes_node);
const auto axes_tensor = ov::Tensor(element::i64, ov::Shape{axes.size()}, axes.data());
const std::unordered_map<size_t, ov::Tensor> constant_data = {{1, axes_tensor}};
std::ostringstream os;
os << "[cpu]squeeze: the shape of input data ";
os << "(";
for (size_t i = 0; i < input_shapes[0].size(); i++) {
os << input_shapes[0][i];
if (i < input_shapes[0].size() - 1) {
os << ".";
}
}
os << ")";
os << " conflicts with the squeeze pattern ";
os << "(";
for (size_t i = 0; i < axes.size(); i++) {
os << axes[i];
if (i < axes.size() - 1) {
os << ".";
}
}
os << ")";
OV_EXPECT_THROW(unit_test::cpu_test_shape_infer(op.get(), input_shapes, output_shapes, constant_data),
ov::Exception,
HasSubstr(os.str()));
unit_test::cpu_test_shape_infer(op.get(), input_shapes, output_shapes, constant_data);
}
INSTANTIATE_TEST_SUITE_P(
CpuShapeInfer,
SqueezeCpuShapeInferenceThrowExceptionTest,
Values(make_tuple(unit_test::ShapeVector{{1, 2, 3, 1}, {1}}, std::vector<int64_t>{1}, StaticShape({})),
make_tuple(unit_test::ShapeVector{{1, 2, 3, 8}, {1}}, std::vector<int64_t>{2}, StaticShape({})),
make_tuple(unit_test::ShapeVector{{1, 2, 3, 8}, {2}}, std::vector<int64_t>{1, 2}, StaticShape({})),
make_tuple(unit_test::ShapeVector{{1, 2, 3, 8}, {1}}, std::vector<int64_t>{-1}, StaticShape({})),
make_tuple(unit_test::ShapeVector{{1, 2, 3, 8}, {2}}, std::vector<int64_t>{-1, -1}, StaticShape({})),
make_tuple(unit_test::ShapeVector{{1, 2, 3, 8}, {2}}, std::vector<int64_t>{-1, -2}, StaticShape({}))),
SqueezeCpuShapeInferenceThrowExceptionTest::getTestCaseName);
CpuShapeInfer_Squeeze_non_squeezable,
SqueezeCpuShapeInferenceTestNonSqueeezable,
Values(
make_tuple(unit_test::ShapeVector{{1, 2, 3, 1}, {1}}, std::vector<int64_t>{1}, StaticShape({1, 2, 3, 1})),
make_tuple(unit_test::ShapeVector{{1, 2, 3, 1}, {2}}, std::vector<int64_t>{0, 1}, StaticShape({2, 3, 1})),
make_tuple(unit_test::ShapeVector{{1, 2, 3, 8}, {1}}, std::vector<int64_t>{2}, StaticShape({1, 2, 3, 8})),
make_tuple(unit_test::ShapeVector{{1, 2, 3, 8}, {2}}, std::vector<int64_t>{1, 2}, StaticShape({1, 2, 3, 8})),
make_tuple(unit_test::ShapeVector{{1, 2, 3, 8}, {1}}, std::vector<int64_t>{-1}, StaticShape({1, 2, 3, 8})),
make_tuple(unit_test::ShapeVector{{1, 2, 3, 8}, {2}}, std::vector<int64_t>{-1, -1}, StaticShape({1, 2, 3, 8})),
make_tuple(unit_test::ShapeVector{{1, 2, 3, 8}, {2}}, std::vector<int64_t>{-1, -2}, StaticShape({1, 2, 3, 8}))),
SqueezeCpuShapeInferenceTest::getTestCaseName);
} // namespace cpu_shape_infer
} // namespace unit_test
} // namespace intel_cpu
} // namespace ov
} // namespace cpu_shape_infer
} // namespace unit_test
} // namespace intel_cpu
} // namespace ov

5
src/plugins/intel_cpu/tests/unit/shape_inference_test/squeeze_shape_inference_test.cpp
View File

@ -2,13 +2,14 @@
// SPDX-License-Identifier: Apache-2.0
//
#include "squeeze_shape_inference.hpp"
#include <gmock/gmock.h>
#include "common_test_utils/test_assertions.hpp"
#include "openvino/op/constant.hpp"
#include "openvino/op/parameter.hpp"
#include "openvino/op/squeeze.hpp"
#include "squeeze_shape_inference.hpp"
#include "utils.hpp"
using namespace ov;
@ -69,6 +70,7 @@ protected:
INSTANTIATE_TEST_SUITE_P(1d_shapes,
SqueezeStaticShapeInferenceTest,
Values(make_tuple(ShapeVector{{1}, {1}}, std::vector<int64_t>{-1}, StaticShape({})),
make_tuple(ShapeVector{{6}, {1}}, std::vector<int64_t>{-1}, StaticShape({6})),
make_tuple(ShapeVector{{1}, {1}}, std::vector<int64_t>{0}, StaticShape({}))),
PrintToStringParamName());
@ -77,6 +79,7 @@ INSTANTIATE_TEST_SUITE_P(
SqueezeStaticShapeInferenceTest,
Values(make_tuple(ShapeVector{{1, 2, 3, 1}, {2}}, std::vector<int64_t>{0, 3}, StaticShape({2, 3})),
make_tuple(ShapeVector{{2, 1, 1, 4}, {2}}, std::vector<int64_t>{2, 1}, StaticShape({2, 4})),
make_tuple(ShapeVector{{2, 1, 1, 4, 1}, {2}}, std::vector<int64_t>{0, 1, -2, -1}, StaticShape({2, 1, 4})),
make_tuple(ShapeVector{{1, 3, 1, 2, 1}, {3}}, std::vector<int64_t>{0, 2, 4}, StaticShape({3, 2})),
make_tuple(ShapeVector{{1, 3, 1, 2, 1}, {3}}, std::vector<int64_t>{4, 2, 0}, StaticShape({3, 2})),
make_tuple(ShapeVector{{1, 3, 1, 2, 1}, {3}}, std::vector<int64_t>{2, 0, 4}, StaticShape({3, 2})),

31
src/plugins/intel_gpu/include/intel_gpu/op/indirect_sdpa.hpp
View File

@ -19,36 +19,7 @@ public:
IndirectSDPA() = default;
IndirectSDPA(const ov::Output<Node>& Q,
const ov::Output<Node>& K,
const ov::Output<Node>& V,
const ov::Output<Node>& beam_table,
const bool is_causal,
const int64_t indirect_axis,
const std::vector<int64_t>& order_q,
const std::vector<int64_t>& order_k,
const std::vector<int64_t>& order_v,
const std::vector<int64_t>& order_out,
const ov::element::Type output_type = ov::element::undefined);
IndirectSDPA(const ov::Output<Node>& Q,
const ov::Output<Node>& K,
const ov::Output<Node>& V,
const ov::Output<Node>& attn_mask,
const ov::Output<Node>& beam_table,
const bool is_causal,
const int64_t indirect_axis,
const std::vector<int64_t>& order_q,
const std::vector<int64_t>& order_k,
const std::vector<int64_t>& order_v,
const std::vector<int64_t>& order_out,
const ov::element::Type output_type = ov::element::undefined);
IndirectSDPA(const ov::Output<Node>& Q,
const ov::Output<Node>& K,
const ov::Output<Node>& V,
const ov::Output<Node>& attn_mask,
const ov::Output<Node>& scale,
IndirectSDPA(const OutputVector& data_inputs,
const ov::Output<Node>& beam_table,
const bool is_causal,
const int64_t indirect_axis,

31
src/plugins/intel_gpu/include/intel_gpu/op/sdpa.hpp
View File

@ -19,37 +19,12 @@ public:
SDPA() = default;
SDPA(const ov::Output<Node>& Q,
const ov::Output<Node>& K,
const ov::Output<Node>& V,
SDPA(const OutputVector& inputs,
const bool is_causal,
const std::vector<int64_t>& order_q,
const std::vector<int64_t>& order_k,
const std::vector<int64_t>& order_v,
const std::vector<int64_t>& order_out,
const bool is_causal,
const ov::element::Type output_type = ov::element::undefined);
SDPA(const ov::Output<Node>& Q,
const ov::Output<Node>& K,
const ov::Output<Node>& V,
const ov::Output<Node>& attn_mask,
const std::vector<int64_t>& order_q,
const std::vector<int64_t>& order_k,
const std::vector<int64_t>& order_v,
const std::vector<int64_t>& order_out,
const bool is_causal,
const ov::element::Type output_type = ov::element::undefined);
SDPA(const ov::Output<Node>& Q,
const ov::Output<Node>& K,
const ov::Output<Node>& V,
const ov::Output<Node>& attn_mask,
const ov::Output<Node>& scale,
const std::vector<int64_t>& order_q,
const std::vector<int64_t>& order_k,
const std::vector<int64_t>& order_v,
const std::vector<int64_t>& order_out,
const bool is_causal,
const ov::element::Type output_type = ov::element::undefined);
bool visit_attributes(ov::AttributeVisitor &visitor) override;
@ -73,11 +48,11 @@ public:
}
protected:
bool m_is_causal;
std::vector<int64_t> m_order_q;
std::vector<int64_t> m_order_k;
std::vector<int64_t> m_order_v;
std::vector<int64_t> m_order_out;
bool m_is_causal;
ov::element::Type m_output_type;
};

1
src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp
View File

@ -286,3 +286,4 @@ REGISTER_FACTORY(internal, Convolution);
REGISTER_FACTORY(internal, Placeholder);
REGISTER_FACTORY(internal, SDPA);
REGISTER_FACTORY(internal, IndirectSDPA);
REGISTER_FACTORY(internal, RoPE);

92
src/plugins/intel_gpu/include/intel_gpu/primitives/rope.hpp Normal file
View File

@ -0,0 +1,92 @@
// Copyright (C) 2023 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#pragma once
#include "primitive.hpp"
#include "ov_ops/rotary_positional_embeddings.hpp"
namespace cldnn {
using RoPE = ov::op::internal::RoPE;
/// @brief Rotary Position Embedding primitive
struct rope : public primitive_base<rope> {
CLDNN_DECLARE_PRIMITIVE(rope);
rope() : primitive_base("", {}) {}
/// @brief Constructs rope primitive
/// @param id This primitive id
/// @param inputs Inputs primitive ids
/// @param config Specific RoPE config
rope(const primitive_id& id,
const std::vector<input_info>& inputs,
const RoPE::Config& config,
const padding& output_padding = padding())
: primitive_base(id, inputs, {output_padding}),
config(config) {}
RoPE::Config config;
size_t hash() const override {
size_t seed = primitive::hash();
seed = hash_combine(seed, config.gather_position_arg_id);
seed = hash_combine(seed, config.head_cnt);
seed = hash_combine(seed, config.head_size);
seed = hash_combine(seed, config.input_trans0213);
seed = hash_combine(seed, config.is_chatglm);
seed = hash_combine(seed, config.is_interleaved);
seed = hash_combine(seed, config.is_qwen);
seed = hash_combine(seed, config.rotary_ndims);
seed = hash_combine(seed, config.slice_start);
seed = hash_combine(seed, config.slice_stop);
return seed;
}
bool operator==(const primitive& rhs) const override {
if (!compare_common_params(rhs))
return false;
auto rhs_casted = downcast<const rope>(rhs);
return config.gather_position_arg_id == rhs_casted.config.gather_position_arg_id &&
config.head_cnt == rhs_casted.config.head_cnt &&
config.head_size == rhs_casted.config.head_size &&
config.input_trans0213 == rhs_casted.config.input_trans0213 &&
config.is_chatglm == rhs_casted.config.is_chatglm &&
config.is_interleaved == rhs_casted.config.is_interleaved &&
config.is_qwen == rhs_casted.config.is_qwen &&
config.rotary_ndims == rhs_casted.config.rotary_ndims &&
config.slice_start == rhs_casted.config.slice_start &&
config.slice_stop == rhs_casted.config.slice_stop;
}
void save(BinaryOutputBuffer& ob) const override {
primitive_base<rope>::save(ob);
ob << config.gather_position_arg_id;
ob << config.head_cnt;
ob << config.head_size;
ob << config.input_trans0213;
ob << config.is_chatglm;
ob << config.is_interleaved;
ob << config.is_qwen;
ob << config.rotary_ndims;
ob << config.slice_start;
ob << config.slice_stop;
}
void load(BinaryInputBuffer& ib) override {
primitive_base<rope>::load(ib);
ib >> config.gather_position_arg_id;
ib >> config.head_cnt;
ib >> config.head_size;
ib >> config.input_trans0213;
ib >> config.is_chatglm;
ib >> config.is_interleaved;
ib >> config.is_qwen;
ib >> config.rotary_ndims;
ib >> config.slice_start;
ib >> config.slice_stop;
}
};
} // namespace cldnn

1
src/plugins/intel_gpu/src/graph/impls/ocl/register.cpp
View File

@ -94,6 +94,7 @@ void register_implementations() {
REGISTER_OCL(unique_count);
REGISTER_OCL(unique_gather);
REGISTER_OCL(scaled_dot_product_attention);
REGISTER_OCL(rope);
}
} // namespace ocl

2
src/plugins/intel_gpu/src/graph/impls/ocl/register.hpp
View File

@ -75,6 +75,7 @@
#include "intel_gpu/primitives/unique.hpp"
#include "intel_gpu/primitives/kv_cache.hpp"
#include "intel_gpu/primitives/scaled_dot_product_attention.hpp"
#include "intel_gpu/primitives/rope.hpp"
namespace cldnn {
namespace ocl {
@ -174,6 +175,7 @@ REGISTER_OCL(eye);
REGISTER_OCL(unique_count);
REGISTER_OCL(unique_gather);
REGISTER_OCL(scaled_dot_product_attention);
REGISTER_OCL(rope);
#undef REGISTER_OCL

88
src/plugins/intel_gpu/src/graph/impls/ocl/rope.cpp Normal file
View File

@ -0,0 +1,88 @@
// Copyright (C) 2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "primitive_base.hpp"
#include "rope_inst.h"
#include "rope/rope_kernel_selector.h"
#include "rope/rope_kernel_ref.h"
namespace cldnn {
namespace ocl {
struct rope_impl : typed_primitive_impl_ocl<rope> {
using parent = typed_primitive_impl_ocl<rope>;
using parent::parent;
using kernel_selector_t = kernel_selector::rope_kernel_selector;
using kernel_params_t = kernel_selector::rope_params;
DECLARE_OBJECT_TYPE_SERIALIZATION(cldnn::ocl::rope_impl);
std::unique_ptr<primitive_impl> clone() const override {
return make_unique<rope_impl>(*this);
}
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);
}
}
static kernel_params_t get_kernel_params(const kernel_impl_params& impl_param, bool is_shape_agnostic = false) {
const auto& primitive = impl_param.typed_desc<rope>();
auto params = get_default_params<kernel_selector::rope_params>(impl_param, is_shape_agnostic);
params.head_cnt = primitive->config.head_cnt;
params.head_size = primitive->config.head_size;
params.rotary_ndims = primitive->config.rotary_ndims;
params.slice_start = primitive->config.slice_start;
params.slice_stop = primitive->config.slice_stop;
params.axis = primitive->config.is_qwen || primitive->config.is_chatglm ? 2 : 3;
params.num_of_inputs = primitive->config.is_chatglm || primitive->config.is_interleaved ? 2 : 3;
params.is_qwen = primitive->config.is_qwen;
params.is_chatglm = primitive->config.is_chatglm;
for (size_t i = 1; i < impl_param.input_layouts.size(); ++i) {
params.inputs.push_back(convert_data_tensor(impl_param.get_input_layout(i)));
}
return params;
}
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);
}
};
namespace detail {
attach_rope_impl::attach_rope_impl() {
auto types = {
data_types::f32,
data_types::f16
};
auto formats = {
format::bfyx
};
implementation_map<rope>::add(impl_types::ocl,
shape_types::any,
typed_primitive_impl_ocl<rope>::create<rope_impl>,
types,
formats);
}
} // namespace detail
} // namespace ocl
} // namespace cldnn
BIND_BINARY_BUFFER_WITH_TYPE(cldnn::ocl::rope_impl)
BIND_BINARY_BUFFER_WITH_TYPE(cldnn::rope)

1
src/plugins/intel_gpu/src/graph/impls/ocl/scaled_dot_product_attention.cpp
View File

@ -72,7 +72,6 @@ protected:
}
static size_t get_beam_table_id(std::shared_ptr<const scaled_dot_product_attention> primitive) {
GPU_DEBUG_TRACE << "get_beam_table_id " << primitive->input_size() - 1 << "\n";
return primitive->input_size() - 1;
}

39
src/plugins/intel_gpu/src/graph/include/rope_inst.h Normal file
View File

@ -0,0 +1,39 @@
// Copyright (C) 2023 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#pragma once
#include "intel_gpu/primitives/rope.hpp"
#include "primitive_inst.h"
#include <string>
namespace cldnn {
template <>
struct typed_program_node<rope> : public typed_program_node_base<rope> {
using parent = typed_program_node_base<rope>;
public:
using parent::parent;
program_node& input(size_t idx = 0) const { return get_dependency(idx); }
std::vector<size_t> get_shape_infer_dependencies() const override { return {}; }
};
using rope_node = typed_program_node<rope>;
template <>
class typed_primitive_inst<rope> : public typed_primitive_inst_base<rope> {
using parent = typed_primitive_inst_base<rope>;
using parent::parent;
public:
template<typename ShapeType>
static std::vector<layout> calc_output_layouts(const rope_node& /*node*/, const kernel_impl_params& impl_param);
static layout calc_output_layout(rope_node const& node, kernel_impl_params const& impl_param);
static std::string to_string(rope_node const& node);
};
using rope_inst = typed_primitive_inst<rope>;
} // namespace cldnn

4
src/plugins/intel_gpu/src/graph/primitive_inst.cpp
View File

@ -1034,8 +1034,8 @@ void primitive_inst::do_runtime_in_place_kv_cache() {
const auto& sequence_axis = desc->concat_axis;
const auto& gather_axis = desc->gather_axis;
const auto& prev_batch_size = past_layout.get_shape()[gather_axis];
const auto& beam_size = present_layout.get_shape()[gather_axis];
const auto& prev_batch_size = static_cast<size_t>(past_layout.get_shape()[gather_axis]);
const auto& beam_size = static_cast<size_t>(present_layout.get_shape()[gather_axis]);
if (prev_batch_size != beam_size) {
// If the previous batch size is not same as beam size, need explicit concat
_impl_params->_can_be_optimized = false;

76
src/plugins/intel_gpu/src/graph/rope.cpp Normal file
View File

@ -0,0 +1,76 @@
// Copyright (C) 2023 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "rope_inst.h"
#include "primitive_type_base.h"
#include "json_object.h"
#include <string>
namespace cldnn {
GPU_DEFINE_PRIMITIVE_TYPE_ID(rope)
layout rope_inst::calc_output_layout(rope_node const& node, kernel_impl_params const& impl_param) {
return calc_output_layouts<ov::PartialShape>(node, impl_param)[0];
}
template<typename ShapeType>
std::vector<layout> rope_inst::calc_output_layouts(rope_node const& node, kernel_impl_params const& impl_param) {
auto desc = impl_param.typed_desc<rope>();
const auto& input0_layout = impl_param.get_input_layout(0);
const auto& input0_shape = input0_layout.get<ShapeType>();
auto output_format = input0_layout.format;
auto output_type = desc->output_data_types[0].value_or(input0_layout.data_type);
if (impl_param.has_fused_primitives()) {
output_type = impl_param.get_output_element_type();
}
ShapeType output_shape = input0_shape;
if (desc->config.is_qwen || desc->config.is_chatglm) {
output_shape = { input0_shape[0],
input0_shape[1],
ov::Dimension(desc->config.head_cnt),
ov::Dimension(desc->config.head_size) };
} else {
auto input_slice_size = desc->config.slice_stop - desc->config.slice_start;
if (input_slice_size > 0) {
output_shape[3] = input_slice_size;
}
if (desc->config.input_trans0213 || desc->config.is_interleaved) {
std::swap(output_shape[2], output_shape[1]);
}
}
return { layout(output_shape, output_type, output_format) };
}
template std::vector<layout> rope_inst::calc_output_layouts<ov::PartialShape>(rope_node const& node, const kernel_impl_params& impl_param);
std::string rope_inst::to_string(rope_node const& node) {
auto desc = node.get_primitive();
auto node_info = node.desc_to_json();
std::stringstream primitive_description;
json_composite rope_info;
rope_info.add("gather_position_arg_id", desc->config.gather_position_arg_id);
rope_info.add("head_cnt", desc->config.head_cnt);
rope_info.add("head_size", desc->config.head_size);
rope_info.add("input_trans0213", desc->config.input_trans0213);
rope_info.add("is_chatglm", desc->config.is_chatglm);
rope_info.add("is_interleaved", desc->config.is_interleaved);
rope_info.add("is_qwen", desc->config.is_qwen);
rope_info.add("rotary_ndims", desc->config.rotary_ndims);
rope_info.add("slice_start", desc->config.slice_start);
rope_info.add("slice_stop", desc->config.slice_stop);
node_info->add("rope info", rope_info);
node_info->dump(primitive_description);
return primitive_description.str();
}
} // namespace cldnn

22
src/plugins/intel_gpu/src/graph/scaled_dot_product_attention.cpp
View File

@ -20,7 +20,27 @@ layout scaled_dot_product_attention_inst::calc_output_layout(scaled_dot_product_
kernel_impl_params const& impl_param) {
auto desc = impl_param.typed_desc<scaled_dot_product_attention>();
return impl_param.get_input_layout(0);
auto transpose_shape = [&desc](const ov::PartialShape& shape, const std::vector<int64_t>& order) {
if (desc->input_q_transpose_order.empty())
return shape;
auto shape_transposed = ov::PartialShape(shape);
auto rank_diff = shape.size() - order.size();
for (size_t i = 0; i < order.size(); i++) {
size_t idx = static_cast<size_t>(order[i]);
shape_transposed[i + rank_diff] = shape[idx + rank_diff];
}
return shape_transposed;
};
auto input0_layout = impl_param.get_input_layout(0);
auto default_out_dt = data_type_traits::is_floating_point(input0_layout.data_type) ? input0_layout.data_type : data_types::f32;
auto output_type = desc->output_data_types[0].value_or(default_out_dt);
auto output_format = input0_layout.format;
auto output_shape = transpose_shape(input0_layout.get_partial_shape(), desc->input_q_transpose_order); // output shape matches with Q input shape
return { layout{output_shape, output_type, output_format, desc->output_paddings[0]} };
}
template<typename ShapeType>

86
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/rope_ref.cl Normal file
View File

@ -0,0 +1,86 @@
// Copyright (C) 2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "include/fetch_utils.cl"
#ifdef CHATGLM
KERNEL(rope_ref)(
OPTIONAL_SHAPE_INFO_ARG
const __global INPUT0_TYPE* input,
const __global INPUT1_TYPE* cos_sin,
__global OUTPUT_TYPE* output)
{
const uint p = get_global_id(0);
const uint b = get_global_id(1);
const uint h = get_global_id(2) % HEAD_COUNT;
const uint rf = get_global_id(2) / HEAD_COUNT;
uint r = rf < HALF_ROTARY_NDIMS ? rf * 2 : 0;
uint f = rf < HEAD_SIZE - ROTARY_NDIMS ? rf : 0;
#ifdef ENABLE_SLICE
uint input_idx = GET_DATA_INDEX(SLICED_INPUT0, p, b, h * HEAD_SIZE, 0);
input_idx += SLICED_FROM_START * (p * INPUT0_FEATURE_NUM + b + 1)
+ SLICED_FROM_END * (p * INPUT0_FEATURE_NUM + b);
#else
uint input_idx = INPUT0_GET_INDEX(p, b, h * HEAD_SIZE, 0);
#endif
uint cos_sin_p = p < INPUT1_BATCH_NUM ? p : 0;
uint cos_sin_b = b < INPUT1_FEATURE_NUM ? b : 0;
uint cos_sin_idx = INPUT1_GET_INDEX(cos_sin_p, cos_sin_b, 0, 0);
uint output_idx = OUTPUT_GET_INDEX(p, b, h, 0);
INPUT1_TYPE cosv = cos_sin[cos_sin_idx + r];
INPUT1_TYPE sinv = cos_sin[cos_sin_idx + r + 1];
INPUT0_TYPE in1 = input[input_idx + r];
INPUT0_TYPE in2 = input[input_idx + r + 1];
output[output_idx + r] = cosv * in1 - sinv * in2;
output[output_idx + r + 1] = sinv * in1 + cosv * in2;
#ifdef ENABLE_IO_COPY
output[output_idx + ROTARY_NDIMS + f] = input[input_idx + ROTARY_NDIMS + f];
#endif
}
#endif
#ifdef QWEN
KERNEL(rope_ref)(
OPTIONAL_SHAPE_INFO_ARG
const __global INPUT0_TYPE* input,
const __global INPUT1_TYPE* cos,
const __global INPUT1_TYPE* sin,
__global OUTPUT_TYPE* output)
{
const uint b = get_global_id(0);
const uint p = get_global_id(1);
const uint h = get_global_id(2) / HALF_ROTARY_NDIMS;
const uint r = get_global_id(2) % HALF_ROTARY_NDIMS;
#ifdef ENABLE_SLICE
uint input_idx = GET_DATA_INDEX(SLICED_INPUT0, b, p, h * HEAD_SIZE, 0);
input_idx += SLICED_FROM_START * (b * INPUT0_FEATURE_NUM + p + 1)
+ SLICED_FROM_END * (b * INPUT0_FEATURE_NUM + p);
#else
uint input_idx = INPUT0_GET_INDEX(b, p, h * HEAD_SIZE, 0);
#endif
uint cos_sin_b = b < INPUT1_BATCH_NUM ? b : 0;
uint cos_sin_p = p + INPUT1_FEATURE_NUM - INPUT0_FEATURE_NUM < INPUT1_FEATURE_NUM ? p + INPUT1_FEATURE_NUM - INPUT0_FEATURE_NUM : 0;
uint cos_sin_h = h < INPUT1_SIZE_Y ? h : 0;
uint cos_sin_idx = INPUT1_GET_INDEX(cos_sin_b, cos_sin_p, cos_sin_h, 0);
uint output_idx = OUTPUT_GET_INDEX(b, p, h, 0);
INPUT0_TYPE in1 = input[input_idx + r];
INPUT0_TYPE in2 = input[input_idx + HALF_ROTARY_NDIMS + r];
output[output_idx + r] = cos[cos_sin_idx + r] * in1 - sin[cos_sin_idx + r] * in2;
output[output_idx + HALF_ROTARY_NDIMS + r] = cos[cos_sin_idx + HALF_ROTARY_NDIMS + r] * in2 +
sin[cos_sin_idx + HALF_ROTARY_NDIMS + r] * in1;
}
#endif

38
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/sdpa_opt.cl
View File

@ -20,16 +20,16 @@
inline uint FUNC(get_input0_index_nt)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w, uint z, uint y, uint x) {
#if INPUT0_SIMPLE
return GET_DATA_INDEX_6D_SAFE(INPUT0, b, f, w, z, y, x);
return GET_DATA_INDEX_6D(INPUT0, b, f, w, z, y, x);
#else
#if INPUT0_DIMS == 4
return INPUT0_GET_INDEX_SAFE(b, f, y, x);
return INPUT0_GET_INDEX(b, f, y, x);
#elif INPUT0_DIMS == 5
return INPUT0_GET_INDEX_SAFE(b, f, z, y, x);
return INPUT0_GET_INDEX(b, f, z, y, x);
#elif INPUT0_DIMS == 6
return INPUT0_GET_INDEX_SAFE(b, f, w, z, y, x);
return INPUT0_GET_INDEX(b, f, w, z, y, x);
#else
# error sdpa_ref.cl : Unsupported input 0 format
# error sdpa_opt.cl : Unsupported input 0 format
#endif
#endif
}
@ -47,16 +47,16 @@ inline uint FUNC(get_input1_index_nt)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, ui
DO_BROADCAST_KEY_VALUE;
#endif
#if INPUT1_SIMPLE
return GET_DATA_INDEX_6D_SAFE(INPUT1, b, f, w, z, y, x);
return GET_DATA_INDEX_6D(INPUT1, b, f, w, z, y, x);
#else
#if INPUT1_DIMS == 4
return INPUT1_GET_INDEX_SAFE(b, f, y, x);
return INPUT1_GET_INDEX(b, f, y, x);
#elif INPUT1_DIMS == 5
return INPUT1_GET_INDEX_SAFE(b, f, z, y, x);
return INPUT1_GET_INDEX(b, f, z, y, x);
#elif INPUT1_DIMS == 6
return INPUT1_GET_INDEX_SAFE(b, f, w, z, y, x);
return INPUT1_GET_INDEX(b, f, w, z, y, x);
#else
# error sdpa_ref.cl : Unsupported input 1 format
# error sdpa_opt.cl : Unsupported input 1 format
#endif
#endif
}
@ -77,13 +77,13 @@ inline uint FUNC(get_input2_index_nt)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, ui
return GET_DATA_INDEX_6D_SAFE(INPUT2, b, f, w, z, y, x);
#else
#if INPUT2_DIMS == 4
return INPUT2_GET_INDEX_SAFE(b, f, y, x);
return INPUT2_GET_INDEX(b, f, y, x);
#elif INPUT2_DIMS == 5
return INPUT2_GET_INDEX_SAFE(b, f, z, y, x);
return INPUT2_GET_INDEX(b, f, z, y, x);
#elif INPUT2_DIMS == 6
return INPUT2_GET_INDEX_SAFE(b, f, w, z, y, x);
return INPUT2_GET_INDEX(b, f, w, z, y, x);
#else
# error sdpa_ref.cl : Unsupported input 1 format
# error sdpa_opt.cl : Unsupported input 1 format
#endif
#endif
}
@ -101,7 +101,7 @@ inline uint FUNC(get_bt_index_nt)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w
#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
# error sdpa_opt.cl : Unsupported beam table format
#endif
}
@ -466,8 +466,8 @@ KERNEL(sdpa_opt)(
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);
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)];
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);
@ -715,7 +715,7 @@ KERNEL(sdpa_opt)(
#endif
#if INPUT0_TYPE_SIZE == 2
/* Adding this clamp improves performance for some reason */
acc[i] = SOFTMAX_ACCUMULATOR_MIN_FUNC(SOFTMAX_ACCUMULATOR_MAX_FUNC(acc[i], INPUT0_VAL_MIN), INPUT0_VAL_MAX);
acc[i] = INPUT0_MIN_FUNC(INPUT0_MAX_FUNC(acc[i], INPUT0_VAL_MIN), INPUT0_VAL_MAX);
#endif
if (seq_len + i >= partition_seq_len) {
acc[i] = INPUT0_VAL_MIN;
@ -749,7 +749,7 @@ KERNEL(sdpa_opt)(
#define QK_ITERS_END QK_MAX_NUMS_PER_SG
#endif
OUTPUT_TYPE qk_max[QK_MAX_NUMS_PER_SG];
SOFTMAX_ACCUMULATOR_TYPE qk_max[QK_MAX_NUMS_PER_SG];
for (uint i = 0; i < QK_MAX_NUMS_PER_SG; i++)
qk_max[i] = SOFTMAX_ACCUMULATOR_VAL_MIN;

1
src/plugins/intel_gpu/src/kernel_selector/common_types.h
View File

@ -97,6 +97,7 @@ enum class KernelType {
UNIQUE_GATHER,
RMS,
SWIGLU,
ROPE
};
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

114
src/plugins/intel_gpu/src/kernel_selector/kernels/rope/rope_kernel_base.cpp Normal file
View File

@ -0,0 +1,114 @@
// Copyright (C) 2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "rope_kernel_base.h"
#include "kernel_selector_utils.h"
namespace kernel_selector {
bool RoPEKernelBase::Validate(const Params& p) const {
return KernelBaseOpenCL::Validate(p);
}
JitConstants RoPEKernelBase::GetJitConstants(const rope_params& params, RoPEKernelBase::DispatchData) const {
JitConstants jit = MakeBaseParamsJitConstants(params);
jit.AddConstant(MakeJitConstant("HEAD_SIZE", params.head_size));
jit.AddConstant(MakeJitConstant("ROTARY_NDIMS", params.rotary_ndims));
jit.AddConstant(MakeJitConstant("HALF_ROTARY_NDIMS", params.rotary_ndims / 2));
jit.AddConstant(MakeJitConstant("HEAD_COUNT", params.head_cnt));
if (params.head_size > params.rotary_ndims) {
jit.AddConstant(MakeJitConstant("ENABLE_IO_COPY", true));
}
if (params.slice_stop - params.slice_start > 0) {
jit.AddConstant(MakeJitConstant("ENABLE_SLICE", true));
auto f = toCodeString(params.inputs[0].Feature(), 1);
auto x = toCodeString(params.inputs[0].X(), 2);
auto y = toCodeString(params.inputs[0].Y(), 3);
auto sliced_y = toCodeString(params.slice_stop - params.slice_start);
jit.AddConstant(MakeJitConstant("SLICED_INPUT0_X_PITCH", 1));
jit.AddConstant(MakeJitConstant("SLICED_INPUT0_Y_PITCH", x));
jit.AddConstant(MakeJitConstant("SLICED_INPUT0_FEATURE_PITCH", x + "*" + sliced_y));
jit.AddConstant(MakeJitConstant("SLICED_INPUT0_BATCH_PITCH", x + "*" + sliced_y + "*" + f));
jit.AddConstant(MakeJitConstant("SLICED_INPUT0_OFFSET", 0));
jit.AddConstant(MakeJitConstant("SLICED_FROM_START", toCodeString(params.slice_start)));
jit.AddConstant(MakeJitConstant("SLICED_FROM_END", "(" + y + "-" + toCodeString(params.slice_stop) + ")"));
}
if (params.is_qwen) {
jit.AddConstant(MakeJitConstant("QWEN", true));
} else if (params.is_chatglm) {
jit.AddConstant(MakeJitConstant("CHATGLM", true));
}
return jit;
}
RoPEKernelBase::DispatchData RoPEKernelBase::SetDefault(const rope_params& params) const {
DispatchData dispatchData;
const auto& input = params.inputs[0];
const auto& output = params.outputs[0];
std::vector<std::vector<Tensor::DataChannelName>> dims_by_gws = {{ Tensor::DataChannelName::BATCH },
{ Tensor::DataChannelName::FEATURE },
{ Tensor::DataChannelName::Y, Tensor::DataChannelName::X }};
dispatchData.gws = {input.Batch().v,
input.Feature().v,
params.head_cnt * std::max(params.rotary_ndims / 2ul, params.head_size - params.rotary_ndims)};
dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, params.engineInfo, input.GetLayout(), output.GetLayout(), dims_by_gws);
return dispatchData;
}
void RoPEKernelBase::GetUpdateDispatchDataFunc(KernelData& kd) const {
kd.update_dispatch_data_func = [this](const Params& params, KernelData& kd) {
const auto& prim_params = static_cast<const rope_params&>(params);
auto dispatchData = SetDefault(prim_params);
OPENVINO_ASSERT(kd.kernels.size() == 1, "[GPU] Invalid kernels size for update dispatch data func");
kd.kernels[0].params.workGroups.global = dispatchData.gws;
kd.kernels[0].params.workGroups.local = dispatchData.lws;
kd.kernels[0].skip_execution = KernelData::SkipKernelExecution(prim_params);
};
}
KernelsData RoPEKernelBase::GetCommonKernelsData(const Params& params) const {
assert(params.GetType() == KernelType::ROPE);
if (!Validate(params))
return {};
const rope_params& orgParams = static_cast<const rope_params&>(params);
auto dispatchData = SetDefault(orgParams);
KernelData kd = KernelData::Default<rope_params>(params);
auto cldnn_jit = GetJitConstants(orgParams, dispatchData);
auto entry_point = GetEntryPoint(kernelName, orgParams.layerID, params);
auto jit = CreateJit(kernelName, cldnn_jit, entry_point);
GetUpdateDispatchDataFunc(kd);
auto& kernel = kd.kernels[0];
FillCLKernelData(kernel,
dispatchData,
params.engineInfo,
kernelName,
jit,
entry_point,
EXE_MODE_DEFAULT,
false,
false,
static_cast<int>(orgParams.num_of_inputs),
GetFusedPrimitiveInputsCount(params),
1,
orgParams.outputs[0].is_dynamic());
return {kd};
}
} // namespace kernel_selector

45
src/plugins/intel_gpu/src/kernel_selector/kernels/rope/rope_kernel_base.h Normal file
View File

@ -0,0 +1,45 @@
// Copyright (C) 2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#pragma once
#include "kernel_base_opencl.h"
namespace kernel_selector {
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
// rope_params
///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
struct rope_params : public base_params {
rope_params() : base_params(KernelType::ROPE) {}
size_t head_cnt;
size_t head_size;
size_t rotary_ndims;
size_t slice_start;
size_t slice_stop;
size_t axis;
size_t num_of_inputs;
bool is_qwen;
bool is_chatglm;
};
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
// RoPEKernelBase
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
class RoPEKernelBase : public KernelBaseOpenCL {
public:
using KernelBaseOpenCL::KernelBaseOpenCL;
virtual ~RoPEKernelBase() {}
struct DispatchData : public CommonDispatchData {};
protected:
bool Validate(const Params&) const override;
virtual JitConstants GetJitConstants(const rope_params& params, DispatchData dispatchData) const;
virtual DispatchData SetDefault(const rope_params& params) const;
KernelsData GetCommonKernelsData(const Params& params) const;
void GetUpdateDispatchDataFunc(KernelData& kd) const override;
};
} // namespace kernel_selector

34
src/plugins/intel_gpu/src/kernel_selector/kernels/rope/rope_kernel_ref.cpp Normal file
View File

@ -0,0 +1,34 @@
// Copyright (C) 2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "rope_kernel_ref.h"
#include "kernel_selector_utils.h"
#include <string>
namespace kernel_selector {
ParamsKey RoPEKernelRef::GetSupportedKey() const {
ParamsKey k;
k.EnableInputDataType(Datatype::F16);
k.EnableInputDataType(Datatype::F32);
k.EnableOutputDataType(Datatype::F16);
k.EnableOutputDataType(Datatype::F32);
k.EnableInputLayout(DataLayout::bfyx);
k.EnableOutputLayout(DataLayout::bfyx);
k.EnableTensorOffset();
k.EnableTensorPitches();
k.EnableBatching();
k.EnableDifferentTypes();
k.EnableDynamicShapesSupport();
return k;
}
KernelsData RoPEKernelRef::GetKernelsData(const Params& params) const {
return GetCommonKernelsData(params);
}
KernelsPriority RoPEKernelRef::GetKernelsPriority(const Params& /*params*/) const {
return FORCE_PRIORITY_9;
}
} // namespace kernel_selector

20
src/plugins/intel_gpu/src/kernel_selector/kernels/rope/rope_kernel_ref.h Normal file
View File

@ -0,0 +1,20 @@
// Copyright (C) 2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#pragma once
#include "rope_kernel_base.h"
namespace kernel_selector {
class RoPEKernelRef : public RoPEKernelBase {
public:
using Parent = RoPEKernelBase;
RoPEKernelRef() : RoPEKernelBase("rope_ref") {}
virtual ~RoPEKernelRef() {}
KernelsData GetKernelsData(const Params& params) const override;
KernelsPriority GetKernelsPriority(const Params& params) const override;
ParamsKey GetSupportedKey() const override;
};
} // namespace kernel_selector

16
src/plugins/intel_gpu/src/kernel_selector/kernels/rope/rope_kernel_selector.cpp Normal file
View File

@ -0,0 +1,16 @@
// Copyright (C) 2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "rope_kernel_selector.h"
#include "rope_kernel_ref.h"
namespace kernel_selector {
rope_kernel_selector::rope_kernel_selector() {
Attach<RoPEKernelRef>();
}
KernelsData rope_kernel_selector::GetBestKernels(const Params& params) const {
return GetNaiveBestKernel(params, KernelType::ROPE);
}
} // namespace kernel_selector

23
src/plugins/intel_gpu/src/kernel_selector/kernels/rope/rope_kernel_selector.h Normal file
View File

@ -0,0 +1,23 @@
// Copyright (C) 2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#pragma once
#include "kernel_selector.h"
namespace kernel_selector {
class rope_kernel_selector : public kernel_selector_base {
public:
static rope_kernel_selector& Instance() {
static rope_kernel_selector instance_;
return instance_;
}
rope_kernel_selector();
virtual ~rope_kernel_selector() {}
KernelsData GetBestKernels(const Params& params) const override;
};
} // namespace kernel_selector

31
src/plugins/intel_gpu/src/kernel_selector/kernels/sdpa/sdpa_kernel_opt.cpp
View File

@ -23,12 +23,15 @@ static size_t get_target_seq_len_block_size() {
return block_size;
}
static size_t get_seq_len_partition_size() {
const size_t seq_len = 256;
return seq_len;
}
static Datatype get_softmax_acc_type() {
return Datatype::F32;
}
static std::string GetKernelName(std::string base_name, KernelsTypes type, bool is_indirect) {
auto kernel_name = base_name;
if (is_indirect)
@ -80,7 +83,7 @@ bool SDPAKernelOpt::Validate(const Params& p) const {
JitConstants SDPAKernelOpt::GetJitConstants(const sdpa_params& params, size_t kernel_idx) const {
auto jit = SDPAKernelBase::GetJitConstants(params);
const auto softmax_acc_dt = params.inputs[0].GetDType();
const auto softmax_acc_dt = get_softmax_acc_type();
jit.Merge(MakeTypeJitConstants(softmax_acc_dt, "SOFTMAX_ACCUMULATOR"));
const auto& config = params.conf;
@ -232,7 +235,7 @@ void SDPAKernelOpt::GetUpdateDispatchDataFunc(KernelData& kd) const {
auto num_of_partitions = CeilDiv(source_seq_len, get_seq_len_partition_size());
auto buf_dt_size = output.ElementSize();
auto buf_dt_size = BytesPerElement(get_softmax_acc_type());
auto buf_elements_count = (num_of_partitions == 1) ? 1 : output.LogicalSize() / head_size * num_of_partitions;
auto buf_size = buf_elements_count * buf_dt_size;
@ -241,26 +244,26 @@ void SDPAKernelOpt::GetUpdateDispatchDataFunc(KernelData& kd) const {
auto tmp_out_size = tmp_out_elements_count * tmp_out_dt_size;
auto dispatch_data1 = SetDefault(prim_params, 0);
kernel_data.kernels[0].params.workGroups.global = dispatch_data1.gws;
kernel_data.kernels[0].params.workGroups.local = dispatch_data1.lws;
kernel_data.kernels[0].skip_execution = target_seq_len > 1;
kernel_data.kernels[KernelsTypes::SINGLE_TOKEN].params.workGroups.global = dispatch_data1.gws;
kernel_data.kernels[KernelsTypes::SINGLE_TOKEN].params.workGroups.local = dispatch_data1.lws;
kernel_data.kernels[KernelsTypes::SINGLE_TOKEN].skip_execution = target_seq_len > 1;
auto dispatch_data2 = SetDefault(prim_params, 1);
kernel_data.kernels[1].params.workGroups.global = dispatch_data2.gws;
kernel_data.kernels[1].params.workGroups.local = dispatch_data2.lws;
kernel_data.kernels[1].skip_execution = target_seq_len == 1;
kernel_data.kernels[KernelsTypes::MULTI_TOKENS].params.workGroups.global = dispatch_data2.gws;
kernel_data.kernels[KernelsTypes::MULTI_TOKENS].params.workGroups.local = dispatch_data2.lws;
kernel_data.kernels[KernelsTypes::MULTI_TOKENS].skip_execution = target_seq_len == 1;
ScalarDescriptor num_of_partitions_scalar;
num_of_partitions_scalar.t = ScalarDescriptor::Types::UINT32;
num_of_partitions_scalar.v.u32 = static_cast<uint32_t>(num_of_partitions);
auto dispatch_data3 = SetDefault(prim_params, 2);
kernel_data.kernels[2].params.workGroups.global = dispatch_data3.gws;
kernel_data.kernels[2].params.workGroups.local = dispatch_data3.lws;
kernel_data.kernels[2].skip_execution = num_of_partitions == 1;
kernel_data.kernels[KernelsTypes::FINALIZATION].params.workGroups.global = dispatch_data3.gws;
kernel_data.kernels[KernelsTypes::FINALIZATION].params.workGroups.local = dispatch_data3.lws;
kernel_data.kernels[KernelsTypes::FINALIZATION].skip_execution = num_of_partitions == 1;
kernel_data.kernels[2].params.scalars.clear();
kernel_data.kernels[2].params.scalars.push_back(num_of_partitions_scalar);
kernel_data.kernels[KernelsTypes::FINALIZATION].params.scalars.clear();
kernel_data.kernels[KernelsTypes::FINALIZATION].params.scalars.push_back(num_of_partitions_scalar);
kernel_data.internalBufferSizes.clear();
kernel_data.internalBufferSizes.push_back(buf_size);

1
src/plugins/intel_gpu/src/plugin/graph.cpp
View File

@ -198,6 +198,7 @@ std::shared_ptr<ov::Model> Graph::get_runtime_model(std::vector<cldnn::primitive
{ "quantize", "Quantize" },
{ "region_yolo", "RegionYolo" },
{ "reorder", "Reorder" },
{ "rope", "RoPE" },
{ "reorg_yolo", "ReorgYolo" },
{ "reshape", "Reshape" },
{ "reverse_sequence", "ReverseSequence" },

37
src/plugins/intel_gpu/src/plugin/ops/rope.cpp Normal file
View File

@ -0,0 +1,37 @@
// Copyright (C) 2023 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "ov_ops/rotary_positional_embeddings.hpp"
#include "intel_gpu/plugin/program_builder.hpp"
#include "intel_gpu/plugin/common_utils.hpp"
#include "intel_gpu/primitives/rope.hpp"
#include "intel_gpu/primitives/permute.hpp"
namespace ov {
namespace op {
namespace internal {
using RoPE = ov::op::internal::RoPE;
} // namespace internal
} // namespace op
} // namespace ov
namespace ov {
namespace intel_gpu {
static void CreateRoPEOp(ProgramBuilder& p, const std::shared_ptr<op::internal::RoPE>& op) {
validate_inputs_count(op, {3, 4});
auto inputs = p.GetInputInfo(op);
const auto& config = op->get_config();
auto rope = cldnn::rope(layer_type_name_ID(op),
inputs,
config);
p.add_primitive(*op, rope);
}
REGISTER_FACTORY_IMPL(internal, RoPE);
} // namespace intel_gpu
} // namespace ov

54
src/plugins/intel_gpu/src/plugin/transformations/indirect_kv_cache.cpp
View File

@ -190,45 +190,27 @@ IndirectSDPAOpt::IndirectSDPAOpt() {
auto order_out = sdpa->get_output_transpose_order();
auto is_causal = sdpa->get_causal();
std::shared_ptr<op::SDPA> indirect_sdpa;
if (pattern_map.find(sdpa_without_attn_mask_m) != pattern_map.end()) {
indirect_sdpa = std::make_shared<ov::intel_gpu::op::IndirectSDPA>(sdpa->get_input_node_shared_ptr(0),
sdpa->get_input_node_shared_ptr(1),
sdpa->get_input_node_shared_ptr(2),
indirect_kv_cache_0->output(1), // beam table
is_causal,
gather_axis_1,
order_in0,
order_in1,
order_in2,
order_out);
} else if (pattern_map.find(sdpa_with_attn_mask_m) != pattern_map.end()) {
indirect_sdpa = std::make_shared<ov::intel_gpu::op::IndirectSDPA>(sdpa->get_input_node_shared_ptr(0),
sdpa->get_input_node_shared_ptr(1),
sdpa->get_input_node_shared_ptr(2),
sdpa->get_input_node_shared_ptr(3),
indirect_kv_cache_0->output(1), // beam table
is_causal,
gather_axis_1,
order_in0,
order_in1,
order_in2,
order_out);
OutputVector data_inputs;
data_inputs.push_back(sdpa->get_input_node_shared_ptr(0)); // Q
data_inputs.push_back(sdpa->get_input_node_shared_ptr(1)); // K
data_inputs.push_back(sdpa->get_input_node_shared_ptr(2)); // V
if (pattern_map.find(sdpa_with_attn_mask_m) != pattern_map.end()) {
data_inputs.push_back(sdpa->get_input_source_output(3));
} else if (pattern_map.find(sdpa_with_attn_mask_and_scale_m) != pattern_map.end()) {
indirect_sdpa = std::make_shared<ov::intel_gpu::op::IndirectSDPA>(sdpa->get_input_node_shared_ptr(0),
sdpa->get_input_node_shared_ptr(1),
sdpa->get_input_node_shared_ptr(2),
sdpa->get_input_node_shared_ptr(3),
sdpa->get_input_node_shared_ptr(4),
indirect_kv_cache_0->output(1), // beam table
is_causal,
gather_axis_1,
order_in0,
order_in1,
order_in2,
order_out);
data_inputs.push_back(sdpa->get_input_source_output(3));
data_inputs.push_back(sdpa->get_input_source_output(4));
}
auto indirect_sdpa = std::make_shared<ov::intel_gpu::op::IndirectSDPA>(data_inputs,
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());

68
src/plugins/intel_gpu/src/plugin/transformations/op/indirect_sdpa.cpp
View File

@ -9,9 +9,7 @@ namespace ov {
namespace intel_gpu {
namespace op {
IndirectSDPA::IndirectSDPA(const ov::Output<Node>& Q,
const ov::Output<Node>& K,
const ov::Output<Node>& V,
IndirectSDPA::IndirectSDPA(const OutputVector& data_inputs,
const ov::Output<Node>& beam_table,
const bool is_causal,
const int64_t indirect_axis,
@ -20,62 +18,28 @@ IndirectSDPA::IndirectSDPA(const ov::Output<Node>& Q,
const std::vector<int64_t>& order_v,
const std::vector<int64_t>& order_out,
const ov::element::Type output_type)
: ov::intel_gpu::op::SDPA(Q, K, V, order_q, order_k, order_v, order_out, is_causal, output_type)
: ov::intel_gpu::op::SDPA(data_inputs, is_causal, order_q, order_k, order_v, order_out, output_type)
, m_indirect_axis(indirect_axis) {
set_argument(3, beam_table);
validate_and_infer_types();
}
IndirectSDPA::IndirectSDPA(const ov::Output<Node>& Q,
const ov::Output<Node>& K,
const ov::Output<Node>& V,
const ov::Output<Node>& attn_mask,
const ov::Output<Node>& beam_table,
const bool is_causal,
const int64_t indirect_axis,
const std::vector<int64_t>& order_q,
const std::vector<int64_t>& order_k,
const std::vector<int64_t>& order_v,
const std::vector<int64_t>& order_out,
const ov::element::Type output_type)
: ov::intel_gpu::op::SDPA(Q, K, V, attn_mask, order_q, order_k, order_v, order_out, is_causal, output_type)
, m_indirect_axis(indirect_axis) {
set_argument(4, beam_table);
validate_and_infer_types();
}
IndirectSDPA::IndirectSDPA(const ov::Output<Node>& Q,
const ov::Output<Node>& K,
const ov::Output<Node>& V,
const ov::Output<Node>& attn_mask,
const ov::Output<Node>& scale,
const ov::Output<Node>& beam_table,
const bool is_causal,
const int64_t indirect_axis,
const std::vector<int64_t>& order_q,
const std::vector<int64_t>& order_k,
const std::vector<int64_t>& order_v,
const std::vector<int64_t>& order_out,
const ov::element::Type output_type)
: ov::intel_gpu::op::SDPA(Q, K, V, attn_mask, scale, order_q, order_k, order_v, order_out, is_causal, output_type)
, m_indirect_axis(indirect_axis) {
set_argument(5, beam_table);
auto beam_table_idx = data_inputs.size();
set_argument(beam_table_idx, beam_table);
validate_and_infer_types();
}
std::shared_ptr<ov::Node> IndirectSDPA::clone_with_new_inputs(const ov::OutputVector& new_args) const {
check_new_args_count(this, new_args);
if (new_args.size() == 4) {
return std::make_shared<IndirectSDPA>(new_args.at(0), new_args.at(1), new_args.at(2), new_args.at(3),
m_is_causal, m_indirect_axis, m_order_q, m_order_k, m_order_v, m_order_out, m_output_type);
} else if (new_args.size() == 5) {
return std::make_shared<IndirectSDPA>(new_args.at(0), new_args.at(1), new_args.at(2), new_args.at(3), new_args.at(4),
m_is_causal, m_indirect_axis, m_order_q, m_order_k, m_order_v, m_order_out, m_output_type);
} else {
return std::make_shared<IndirectSDPA>(new_args.at(0), new_args.at(1), new_args.at(2), new_args.at(3), new_args.at(4), new_args.at(5),
m_is_causal, m_indirect_axis, m_order_q, m_order_k, m_order_v, m_order_out, m_output_type);
}
// Exclude beam_table input
OutputVector data_inputs(new_args.begin(), new_args.end() - 1);
return std::make_shared<IndirectSDPA>(data_inputs,
new_args.back(),
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() {

72
src/plugins/intel_gpu/src/plugin/transformations/op/sdpa.cpp
View File

@ -14,65 +14,20 @@ namespace ov {
namespace intel_gpu {
namespace op {
SDPA::SDPA(const ov::Output<Node>& Q,
const ov::Output<Node>& K,
const ov::Output<Node>& V,
SDPA::SDPA(const OutputVector& inputs,
const bool is_causal,
const std::vector<int64_t>& order_q,
const std::vector<int64_t>& order_k,
const std::vector<int64_t>& order_v,
const std::vector<int64_t>& order_out,
const bool is_causal,
const ov::element::Type output_type)
: m_order_q(order_q)
: m_is_causal(is_causal)
, m_order_q(order_q)
, m_order_k(order_k)
, m_order_v(order_v)
, m_order_out(order_out)
, m_is_causal(is_causal)
, m_output_type(output_type) {
set_arguments({Q, K, V});
set_causal(is_causal);
validate_and_infer_types();
}
SDPA::SDPA(const ov::Output<Node>& Q,
const ov::Output<Node>& K,
const ov::Output<Node>& V,
const ov::Output<Node>& attn_mask,
const std::vector<int64_t>& order_q,
const std::vector<int64_t>& order_k,
const std::vector<int64_t>& order_v,
const std::vector<int64_t>& order_out,
const bool is_causal,
const ov::element::Type output_type)
: m_order_q(order_q)
, m_order_k(order_k)
, m_order_v(order_v)
, m_order_out(order_out)
, 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();
}
SDPA::SDPA(const ov::Output<Node>& Q,
const ov::Output<Node>& K,
const ov::Output<Node>& V,
const ov::Output<Node>& attn_mask,
const ov::Output<Node>& scale,
const std::vector<int64_t>& order_q,
const std::vector<int64_t>& order_k,
const std::vector<int64_t>& order_v,
const std::vector<int64_t>& order_out,
const bool is_causal,
const ov::element::Type output_type)
: m_order_q(order_q)
, m_order_k(order_k)
, m_order_v(order_v)
, m_order_out(order_out)
, m_is_causal(is_causal)
, m_output_type(output_type) {
set_arguments({Q, K, V, attn_mask, scale});
set_arguments(inputs);
set_causal(is_causal);
validate_and_infer_types();
}
@ -80,16 +35,13 @@ SDPA::SDPA(const ov::Output<Node>& Q,
std::shared_ptr<ov::Node> SDPA::clone_with_new_inputs(const ov::OutputVector& new_args) const {
check_new_args_count(this, new_args);
if (new_args.size() == 3) {
return std::make_shared<SDPA>(new_args.at(0), new_args.at(1), new_args.at(2),
m_order_q, m_order_k, m_order_v, m_order_out, m_is_causal, m_output_type);
} else if (new_args.size() == 4) {
return std::make_shared<SDPA>(new_args.at(0), new_args.at(1), new_args.at(2), new_args.at(3),
m_order_q, m_order_k, m_order_v, m_order_out, m_is_causal, m_output_type);
} else {
return std::make_shared<SDPA>(new_args.at(0), new_args.at(1), new_args.at(2), new_args.at(3), new_args.at(4),
m_order_q, m_order_k, m_order_v, m_order_out, m_is_causal, m_output_type);
}
return std::make_shared<SDPA>(new_args,
m_is_causal,
m_order_q,
m_order_k,
m_order_v,
m_order_out,
m_output_type);
}
void SDPA::validate_and_infer_types() {

20
src/plugins/intel_gpu/src/plugin/transformations/transpose_fusion.cpp
View File

@ -134,18 +134,20 @@ TransposeSDPAMatcher::TransposeSDPAMatcher() {
auto input_k = ov::Output<Node>(pattern_map.at(input_k_m).get_node_shared_ptr(), input_k_output_idx);
auto input_v = ov::Output<Node>(pattern_map.at(input_v_m).get_node_shared_ptr(), input_v_output_idx);
std::shared_ptr<op::SDPA> sdpa_new;
if (pattern_map.find(sdpa_without_attn_mask_m) != pattern_map.end()) {
sdpa_new = std::make_shared<op::SDPA>(input_q, input_k, input_v, order_q, order_k, order_v, order_output, sdpa->get_causal());
} else if (pattern_map.find(sdpa_with_attn_mask_m) != pattern_map.end()) {
auto attn_mask = sdpa->get_input_source_output(3);
sdpa_new = std::make_shared<op::SDPA>(input_q, input_k, input_v, attn_mask, order_q, order_k, order_v, order_output, sdpa->get_causal());
OutputVector inputs;
inputs.push_back(ov::Output<Node>(pattern_map.at(input_q_m).get_node_shared_ptr(), input_q_output_idx));
inputs.push_back(ov::Output<Node>(pattern_map.at(input_k_m).get_node_shared_ptr(), input_k_output_idx));
inputs.push_back(ov::Output<Node>(pattern_map.at(input_v_m).get_node_shared_ptr(), input_v_output_idx));
if (pattern_map.find(sdpa_with_attn_mask_m) != pattern_map.end()) {
inputs.push_back(sdpa->get_input_source_output(3));
} else if (pattern_map.find(sdpa_with_attn_mask_and_scale_m) != pattern_map.end()) {
auto attn_mask = sdpa->get_input_source_output(3);
auto scale = sdpa->get_input_source_output(4);
sdpa_new = std::make_shared<op::SDPA>(input_q, input_k, input_v, attn_mask, scale, order_q, order_k, order_v, order_output, sdpa->get_causal());
inputs.push_back(sdpa->get_input_source_output(3));
inputs.push_back(sdpa->get_input_source_output(4));
}
auto sdpa_new = std::make_shared<op::SDPA>(inputs, sdpa->get_causal(), order_q, order_k, order_v, order_output);
sdpa_new->set_friendly_name(sdpa->get_friendly_name());
ov::copy_runtime_info(m.get_matched_nodes(), sdpa_new);
ov::replace_node(sdpa, sdpa_new);

26
src/plugins/intel_gpu/src/plugin/transformations/unsqueeze_broadcast_reshape_sdpa_fusion.cpp
View File

@ -97,27 +97,25 @@ UnsqueezeBroadcastReshapeSDPAFusion::UnsqueezeBroadcastReshapeSDPAFusion() {
return false;
}
auto input_a = pattern_map.at(input_a_m).get_node_shared_ptr();
auto input_b = pattern_map.at(input_b_m).get_node_shared_ptr();
auto input_c = pattern_map.at(input_c_m).get_node_shared_ptr();
OutputVector data_inputs;
data_inputs.push_back(pattern_map.at(input_a_m).get_node_shared_ptr()); // Q
data_inputs.push_back(pattern_map.at(input_b_m).get_node_shared_ptr()); // K
data_inputs.push_back(pattern_map.at(input_c_m).get_node_shared_ptr()); // V
auto sdpa = std::dynamic_pointer_cast<op::SDPA>(m.get_match_root());
if (pattern_map.find(sdpa_with_attn_mask_m) != pattern_map.end()) {
data_inputs.push_back(sdpa->get_input_source_output(3)); // attn_mask
} else if (pattern_map.find(sdpa_with_attn_mask_and_scale_m) != pattern_map.end()) {
data_inputs.push_back(sdpa->get_input_source_output(3)); // attn_mask
data_inputs.push_back(sdpa->get_input_source_output(4)); // scale
}
auto order_a = sdpa->get_input0_transpose_order();
auto order_b = sdpa->get_input1_transpose_order();
auto order_c = sdpa->get_input2_transpose_order();
auto order_d = sdpa->get_output_transpose_order();
std::shared_ptr<op::SDPA> sdpa_new;
if (pattern_map.find(sdpa_without_attn_mask_m) != pattern_map.end()) {
sdpa_new = std::make_shared<op::SDPA>(input_a, input_b, input_c, order_a, order_b, order_c, order_d, sdpa->get_causal());
} else if (pattern_map.find(sdpa_with_attn_mask_m) != pattern_map.end()) {
auto attn_mask = sdpa->get_input_source_output(3);
sdpa_new = std::make_shared<op::SDPA>(input_a, input_b, input_c, attn_mask, order_a, order_b, order_c, order_d, sdpa->get_causal());
} else if (pattern_map.find(sdpa_with_attn_mask_and_scale_m) != pattern_map.end()) {
auto attn_mask = sdpa->get_input_source_output(3);
auto scale = sdpa->get_input_source_output(4);
sdpa_new = std::make_shared<op::SDPA>(input_a, input_b, input_c, attn_mask, scale, order_a, order_b, order_c, order_d, sdpa->get_causal());
}
auto sdpa_new = std::make_shared<op::SDPA>(data_inputs, sdpa->get_causal(), order_a, order_b, order_c, order_d);
sdpa_new->set_friendly_name(sdpa->get_friendly_name());
ov::copy_runtime_info(m.get_matched_nodes(), sdpa_new);

12
src/plugins/intel_gpu/src/plugin/transformations_pipeline.cpp
View File

@ -61,14 +61,13 @@
#include "plugin/transformations/kv_cache_fusion.hpp"
#include "plugin/transformations/move_fc_reshape_to_weights.hpp"
#include "plugin/transformations/bcast_and_pad_zp_buffers.hpp"
#include "transformations/common_optimizations/rms_fusion.hpp"
#include "plugin/transformations/swiglu_fusion.hpp"
#include "plugin/transformations/transpose_fusion.hpp"
#include "plugin/transformations/indirect_kv_cache.hpp"
#include "plugin/transformations/convert_convolution.hpp"
#include "plugin/transformations/unsqueeze_broadcast_reshape_matmul_fusion.hpp"
#include "transformations/common_optimizations/rms_fusion.hpp"
#include "plugin/transformations/unsqueeze_broadcast_reshape_sdpa_fusion.hpp"
#include "transformations/common_optimizations/rms_fusion.hpp"
#include "transformations/common_optimizations/broadcast_elementwise_fusion.hpp"
#include "transformations/common_optimizations/broadcast_transition.hpp"
#include "transformations/common_optimizations/common_optimizations.hpp"
@ -81,6 +80,7 @@
#include "transformations/common_optimizations/transpose_sinking.hpp"
#include "transformations/common_optimizations/weights_dequantize_to_fake_quantize.hpp"
#include "transformations/common_optimizations/wrap_interpolate_into_transposes.hpp"
#include "transformations/common_optimizations/fuse_rotary_positional_embeddings.hpp"
#include "transformations/control_flow/unroll_tensor_iterator.hpp"
#include "transformations/convert_pooling_to_reduce.hpp"
#include "transformations/convert_precision.hpp"
@ -815,6 +815,14 @@ void TransformationsPipeline::apply(std::shared_ptr<ov::Model> func) {
const size_t zp_pad_size = device_info.supports_immad ? 16 : 32;
manager.register_pass<ov::intel_gpu::BroadcastAndPadZeroPointBuffers>(zp_pad_size);
manager.register_pass<ov::pass::RoPEFusion>();
pass_config->disable<ov::pass::RoPEFusionGPTNEOX>();
pass_config->disable<ov::pass::RoPEFusionGPTJ>();
pass_config->disable<ov::pass::RoPEFusionCosSinPreprocess>();
pass_config->disable<ov::pass::RoPEFusionIOSlicing>();
pass_config->disable<ov::pass::RoPEFusionPreprocess>();
pass_config->disable<ov::pass::RoPEShareCosSin>();
// This is supposed to be the last pass to ensure that we don't have name collisions until
// GPU plugin stops using friendly names for program creation
manager.register_pass<ov::pass::ResolveNameCollisions>(true);

1
src/plugins/intel_gpu/tests/functional/shared_tests_instances/skip_tests_config.cpp
View File

@ -201,6 +201,7 @@ std::vector<std::string> disabledTestPatterns() {
R"(.*smoke_RDFT_5d_last_axis/RDFTLayerTest.Inference/IS=\(10.4.8.2.5\)_modelType=f32_Axes=\(0.1.2.3.4\)_SignalSize=\(\).*)",
// Issue: 136862
R"(.*smoke_ConditionGPUTest_static/StaticConditionLayerGPUTest.CompareWithRefs/IS=\(3.6\)_netPRC=i8_ifCond=PARAM_targetDevice=GPU_.*)",
#if defined(_WIN32)
// by calc abs_threshold with expected value
R"(.*smoke_RemoteTensor/OVRemoteTensorBatched_Test.NV12toBGR_buffer/(num_batch_4|num_batch_2).*)",

22
src/plugins/intel_gpu/tests/functional/shared_tests_instances/subgraph_tests/rotary_pos_emb.cpp Normal file
View File

@ -0,0 +1,22 @@
// Copyright (C) 2018-2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "subgraph_tests/rotary_pos_emb.hpp"
namespace ov {
namespace test {
INSTANTIATE_TEST_SUITE_P(smoke_RoPETestChatGLM,
RoPETestChatGLM,
::testing::Values(ov::test::utils::DEVICE_GPU),
RoPETestChatGLM::getTestCaseName);
INSTANTIATE_TEST_SUITE_P(smoke_RoPETestQwen7b,
RoPETestQwen7b,
::testing::Combine(::testing::Values(true, false),
::testing::Values(ov::test::utils::DEVICE_GPU)),
RoPETestQwen7b::getTestCaseName);
} // namespace test
} // namespace ov

24
src/plugins/intel_gpu/tests/functional/single_layer_tests/dynamic/scaled_dot_product_attention.cpp
View File

@ -25,8 +25,7 @@ typedef std::tuple<ov::element::Type, // netPrecision
std::vector<InputShape>, // shape
bool, // is_causal
bool, // has_attn
bool, // has_scale
std::string // targetDevice
bool // has_scale
> ScaledAttnGPUTestParams;
class ScaledAttnLayerGPUTest : public testing::WithParamInterface<ScaledAttnGPUTestParams>,
@ -48,8 +47,7 @@ std::string ScaledAttnLayerGPUTest::getTestCaseName(const testing::TestParamInfo
bool is_causal;
bool has_attn;
bool has_scale;
std::string targetDevice;
std::tie(inType, inputShapes, is_causal, has_attn, has_scale, targetDevice) = obj.param;
std::tie(inType, inputShapes, is_causal, has_attn, has_scale) = obj.param;
std::ostringstream result;
result << "netPRC=" << inType << "_";
@ -67,7 +65,6 @@ std::string ScaledAttnLayerGPUTest::getTestCaseName(const testing::TestParamInfo
result << "is_causal=" << is_causal << "_";
result << "has_attn=" << has_attn << "_";
result << "has_scale=" << has_scale << "_";
result << "trgDev=" << targetDevice;
return result.str();
}
@ -75,7 +72,10 @@ std::string ScaledAttnLayerGPUTest::getTestCaseName(const testing::TestParamInfo
void ScaledAttnLayerGPUTest::SetUp() {
ov::element::Type inType;
std::vector<InputShape> inputShapes;
std::tie(inType, inputShapes, is_causal, has_attn, has_scale, targetDevice) = this->GetParam();
targetDevice = ov::test::utils::DEVICE_GPU;
std::tie(inType, inputShapes, is_causal, has_attn, has_scale) = this->GetParam();
init_input_shapes(inputShapes);
ov::ParameterVector inputParams;
@ -166,8 +166,7 @@ TEST_P(ScaledAttnLayerGPUTest, CompareWithRefs) {
bool is_causal;
bool has_attn;
bool has_scale;
std::string targetDevice;
std::tie(inType, inputShapes, is_causal, has_attn, has_scale, targetDevice) = this->GetParam();
std::tie(inType, inputShapes, is_causal, has_attn, has_scale) = this->GetParam();
run();
}
@ -190,15 +189,15 @@ const std::vector<std::vector<InputShape>> shapes{
{
// q shape
{ov::test::InputShape{ov::PartialShape{-1, 5, -1, 128},
{ov::Shape{2, 5, 100, 128}, ov::Shape{2, 5, 1, 128}, ov::Shape{2, 5, 384, 128}}}
{ov::Shape{2, 5, 100, 128}, ov::Shape{2, 5, 1, 128}, ov::Shape{2, 5, 387, 128}}}
},
// kv shape
{ov::test::InputShape{ov::PartialShape{-1, 5, -1, 128},
{ov::Shape{2, 5, 100, 128}, ov::Shape{2, 5, 1, 128}, ov::Shape{2, 5, 384, 128}}}
{ov::Shape{2, 5, 100, 128}, ov::Shape{2, 5, 1, 128}, ov::Shape{2, 5, 387, 128}}}
},
// attn shape: [B, 1, -1, L0+L1]
{ov::test::InputShape{ov::PartialShape{-1, 1, -1, -1},
{ov::Shape{1, 1, 100, 100}, ov::Shape{1, 1, 1, 1}, ov::Shape{2, 1, 384, 384}}}
{ov::Shape{1, 1, 100, 100}, ov::Shape{1, 1, 1, 1}, ov::Shape{2, 1, 387, 387}}}
},
},
// heads number of kv is 1, attn mask: [B, H, L1, L0+L1]
@ -222,8 +221,7 @@ const auto params = testing::Combine(testing::Values(ov::element::f16 /*, ov::el
testing::ValuesIn(shapes),
testing::Values(true, false),
testing::Values(true, false),
testing::Values(true, false),
testing::Values(ov::test::utils::DEVICE_GPU));
testing::Values(true, false));
INSTANTIATE_TEST_SUITE_P(smoke_ScaledAttn_GPU,
ScaledAttnLayerGPUTest,

50
src/plugins/intel_gpu/tests/unit/transformations/indirect_kv_cache_test.cpp
View File

@ -22,6 +22,8 @@
#include "intel_gpu/op/indirect_gemm.hpp"
#include "intel_gpu/op/gemm.hpp"
#include "intel_gpu/op/indirect_sdpa.hpp"
#include "intel_gpu/op/sdpa.hpp"
#include "intel_gpu/op/read_value.hpp"
#include "intel_gpu/op/kv_cache.hpp"
@ -132,3 +134,51 @@ TEST_F(TransformationTestsF, IndirectKVCache3) {
comparator.enable(FunctionsComparator::ATTRIBUTES);
}
}
TEST_F(TransformationTestsF, IndirectKVCache4) {
std::vector<int64_t> in0_order = {0, 1, 2, 3};
std::vector<int64_t> in1_order = {0, 2, 1, 3};
std::vector<int64_t> in2_order = {0, 2, 1, 3};
std::vector<int64_t> out_order = {0, 1, 2, 3};
const bool is_causal = true;
{
auto key_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{1, -1, 32, 128}, ov::element::f32, "v0"});
auto value_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{1, -1, 32, 128}, ov::element::f32, "v1"});
auto parameter_key = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{1, -1, 32, 128});
auto parameter_value = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{1, -1, 32, 128});
auto beam_idx = std::make_shared<ov::op::v0::Parameter>(ov::element::i32, ov::PartialShape{1});
auto key_past = std::make_shared<ov::intel_gpu::op::ReadValue>(key_variable);
auto value_past = std::make_shared<ov::intel_gpu::op::ReadValue>(value_variable);
auto axis = std::make_shared<ov::op::v0::Constant>(ov::element::i64, ov::Shape{}, 1);
auto key_gather_past = std::make_shared<ov::op::v8::Gather>(key_past, beam_idx, axis);
auto value_gather_past = std::make_shared<ov::op::v8::Gather>(value_past, beam_idx, axis);
auto key_cache = std::make_shared<ov::intel_gpu::op::KVCache>(key_gather_past, parameter_key, key_variable, 2, ov::element::f32);
auto value_cache = std::make_shared<ov::intel_gpu::op::KVCache>(value_gather_past, parameter_value, value_variable, 2, ov::element::f32);
auto sdpa_q = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{1, 32, -1, 128});
auto inputs = ov::OutputVector{sdpa_q, key_cache, value_cache};
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(inputs, is_causal, in0_order, in1_order, in2_order, out_order);
auto result = std::make_shared<ov::op::v0::Result>(sdpa);
model = std::make_shared<ov::Model>(ov::ResultVector{result}, ov::ParameterVector{parameter_key, parameter_value, beam_idx, sdpa_q});
manager.register_pass<IndirectKVCache>();
}
{
auto indirect_axis = 1;
auto key_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{1, -1, 32, 128}, ov::element::f32, "v0"});
auto value_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{1, -1, 32, 128}, ov::element::f32, "v1"});
auto parameter_key = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{1, -1, 32, 128});
auto parameter_value = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{1, -1, 32, 128});
auto beam_idx = std::make_shared<ov::op::v0::Parameter>(ov::element::i32, ov::PartialShape{1});
auto key_past = std::make_shared<ov::intel_gpu::op::ReadValue>(key_variable);
auto value_past = std::make_shared<ov::intel_gpu::op::ReadValue>(value_variable);
auto key_cache = std::make_shared<ov::intel_gpu::op::KVCache>(key_past, parameter_key, beam_idx, key_variable, 2, indirect_axis, ov::element::f32);
auto value_cache = std::make_shared<ov::intel_gpu::op::KVCache>(value_past, parameter_value, beam_idx, value_variable, 2, indirect_axis, ov::element::f32);
auto sdpa_q = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{1, 32, -1, 128});
auto inputs = ov::OutputVector{sdpa_q, key_cache, value_cache};
auto sdpa = std::make_shared<ov::intel_gpu::op::IndirectSDPA>(inputs, key_cache->output(1), is_causal, indirect_axis, in0_order, in1_order, in2_order, out_order);
auto result = std::make_shared<ov::op::v0::Result>(sdpa);
model_ref = std::make_shared<ov::Model>(ov::ResultVector{result}, ov::ParameterVector{parameter_key, parameter_value, beam_idx, sdpa_q});
comparator.enable(FunctionsComparator::ATTRIBUTES);
}
}

27
src/plugins/intel_gpu/tests/unit/transformations/transpose_sdpa_fusion_test.cpp
View File

@ -25,11 +25,12 @@ namespace test {
namespace intel_gpu {
TEST_F(TransformationTestsF, TranposeSDPAFusion1) {
const bool is_causal = true;
{
auto input_a = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto input_b = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto input_c = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto sdpa = std::make_shared<ov::op::v13::ScaledDotProductAttention>(input_a, input_b, input_c, true);
auto sdpa = std::make_shared<ov::op::v13::ScaledDotProductAttention>(input_a, input_b, input_c, is_causal);
model = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_a, input_b, input_c });
manager.register_pass<TransposeFusion>();
@ -42,21 +43,22 @@ TEST_F(TransformationTestsF, TranposeSDPAFusion1) {
auto input_a = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto input_b = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto input_c = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(input_a, input_b, input_c, order_a, order_b, order_c, order_output, true, ov::element::undefined );
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(ov::OutputVector{input_a, input_b, input_c}, is_causal, order_a, order_b, order_c, order_output, ov::element::undefined );
model_ref = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_a, input_b, input_c });
comparator.enable(FunctionsComparator::ATTRIBUTES);
}
}
TEST_F(TransformationTestsF, TranposeSDPAFusion2) {
TEST_F(TransformationTestsF, TransformationTestsF) {
const bool is_causal = true;
{
auto input_a = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto tranpose_a_const = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{4}, {0, 2, 1, 3});
auto tranpose_a = std::make_shared<ov::op::v1::Transpose>(input_a, tranpose_a_const);
auto input_b = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto input_c = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto sdpa = std::make_shared<ov::op::v13::ScaledDotProductAttention>(tranpose_a, input_b, input_c, true);
auto sdpa = std::make_shared<ov::op::v13::ScaledDotProductAttention>(tranpose_a, input_b, input_c, is_causal);
model = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_a, input_b, input_c });
manager.register_pass<TransposeFusion>();
@ -69,7 +71,7 @@ TEST_F(TransformationTestsF, TranposeSDPAFusion2) {
auto input_a = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto input_b = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto input_c = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(input_a, input_b, input_c, order_a, order_b, order_c, order_output, true, ov::element::undefined);
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(ov::OutputVector{input_a, input_b, input_c}, is_causal, order_a, order_b, order_c, order_output, ov::element::undefined);
model_ref = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_a, input_b, input_c });
comparator.enable(FunctionsComparator::ATTRIBUTES);
@ -77,6 +79,7 @@ TEST_F(TransformationTestsF, TranposeSDPAFusion2) {
}
TEST_F(TransformationTestsF, TranposeSDPAFusion3) {
const bool is_causal = false;
{
auto input_a = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto tranpose_a_const = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{4}, {0, 2, 1, 3});
@ -86,7 +89,7 @@ TEST_F(TransformationTestsF, TranposeSDPAFusion3) {
auto tranpose_b = std::make_shared<ov::op::v1::Transpose>(input_b, tranpose_b_const);
auto input_c = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto sdpa = std::make_shared<ov::op::v13::ScaledDotProductAttention>(tranpose_a, tranpose_b, input_c, false);
auto sdpa = std::make_shared<ov::op::v13::ScaledDotProductAttention>(tranpose_a, tranpose_b, input_c, is_causal);
model = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_a, input_b, input_c });
manager.register_pass<TransposeFusion>();
@ -99,7 +102,7 @@ TEST_F(TransformationTestsF, TranposeSDPAFusion3) {
auto input_a = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto input_b = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto input_c = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(input_a, input_b, input_c, order_a, order_b, order_c, order_output, false, ov::element::undefined);
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(ov::OutputVector{input_a, input_b, input_c}, is_causal, order_a, order_b, order_c, order_output, ov::element::undefined);
model_ref = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_a, input_b, input_c });
comparator.enable(FunctionsComparator::ATTRIBUTES);
@ -107,6 +110,7 @@ TEST_F(TransformationTestsF, TranposeSDPAFusion3) {
}
TEST_F(TransformationTestsF, TranposeSDPAFusion4) {
const bool is_causal = false;
{
auto input_a = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto tranpose_a_const = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{4}, {0, 2, 1, 3});
@ -118,7 +122,7 @@ TEST_F(TransformationTestsF, TranposeSDPAFusion4) {
auto tranpose_c_const = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{4}, {0, 2, 1, 3});
auto tranpose_c = std::make_shared<ov::op::v1::Transpose>(input_c, tranpose_c_const);
auto sdpa = std::make_shared<ov::op::v13::ScaledDotProductAttention>(tranpose_a, tranpose_b, tranpose_c, false);
auto sdpa = std::make_shared<ov::op::v13::ScaledDotProductAttention>(tranpose_a, tranpose_b, tranpose_c, is_causal);
model = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_a, input_b, input_c });
manager.register_pass<TransposeFusion>();
@ -131,7 +135,7 @@ TEST_F(TransformationTestsF, TranposeSDPAFusion4) {
auto input_a = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto input_b = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto input_c = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(input_a, input_b, input_c, order_a, order_b, order_c, order_output, false, ov::element::undefined);
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(ov::OutputVector{input_a, input_b, input_c}, is_causal, order_a, order_b, order_c, order_output, ov::element::undefined);
model_ref = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_a, input_b, input_c });
comparator.enable(FunctionsComparator::ATTRIBUTES);
@ -139,6 +143,7 @@ TEST_F(TransformationTestsF, TranposeSDPAFusion4) {
}
TEST_F(TransformationTestsF, TranposeSDPAFusion5) {
const bool is_causal = false;
{
auto input_a = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto tranpose_a_const = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{4}, {0, 2, 1, 3});
@ -150,7 +155,7 @@ TEST_F(TransformationTestsF, TranposeSDPAFusion5) {
auto tranpose_c_const = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{4}, {3, 2, 1, 0});
auto tranpose_c = std::make_shared<ov::op::v1::Transpose>(input_c, tranpose_c_const);
auto sdpa = std::make_shared<ov::op::v13::ScaledDotProductAttention>(tranpose_a, tranpose_b, tranpose_c, false);
auto sdpa = std::make_shared<ov::op::v13::ScaledDotProductAttention>(tranpose_a, tranpose_b, tranpose_c, is_causal);
model = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_a, input_b, input_c });
manager.register_pass<TransposeFusion>();
@ -166,7 +171,7 @@ TEST_F(TransformationTestsF, TranposeSDPAFusion5) {
auto tranpose_c_const = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{4}, {3, 2, 1, 0});
auto tranpose_c = std::make_shared<ov::op::v1::Transpose>(input_c, tranpose_c_const);
auto sdpa = std::make_shared<ov::op::v13::ScaledDotProductAttention>(tranpose_a, tranpose_b, tranpose_c, false);
auto sdpa = std::make_shared<ov::op::v13::ScaledDotProductAttention>(tranpose_a, tranpose_b, tranpose_c, is_causal);
model_ref = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_a, input_b, input_c });
comparator.enable(FunctionsComparator::ATTRIBUTES);

233
src/plugins/intel_gpu/tests/unit/transformations/unsqueeze_broadcast_reshape_sdpa_fusion_test.cpp Normal file
View File

@ -0,0 +1,233 @@
// Copyright (C) 2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "common_test_utils/ov_test_utils.hpp"
#include "openvino/core/model.hpp"
#include "openvino/pass/manager.hpp"
#include "openvino/op/abs.hpp"
#include "openvino/op/broadcast.hpp"
#include "openvino/op/constant.hpp"
#include "openvino/op/gather.hpp"
#include "openvino/op/reshape.hpp"
#include "openvino/op/unsqueeze.hpp"
#include "intel_gpu/op/sdpa.hpp"
#include "intel_gpu/op/read_value.hpp"
#include "intel_gpu/op/kv_cache.hpp"
#include "plugin/transformations/unsqueeze_broadcast_reshape_sdpa_fusion.hpp"
#include <memory>
using namespace testing;
using namespace ov::intel_gpu;
namespace ov {
namespace test {
namespace intel_gpu {
TEST_F(TransformationTestsF, UnsqueezeBroadReshapeSDPAFusion1) {
std::vector<int64_t> in0_order = {0, 2, 1, 3};
std::vector<int64_t> in1_order = {0, 2, 1, 3};
std::vector<int64_t> in2_order = {0, 2, 1, 3};
std::vector<int64_t> out_order = {0, 1, 2, 3};
std::vector<int64_t> axes_val = {-2};
std::vector<int32_t> target_shape_kv = {1, 1, 1, 16, 1};
std::vector<int64_t> pattern_shape = {0, 0, 32, 32};
const bool is_causal = true;
{
auto input_q = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto key_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{-1, -1, 2, 32}, ov::element::f32, "v0"});
auto value_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{-1, -1, 2, 32}, ov::element::f32, "v1"});
auto key_token_param = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{-1, -1, 2, 32});
auto value_token_param = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{-1, -1, 2, 32});
auto beam_idx = std::make_shared<ov::op::v0::Parameter>(ov::element::i32, ov::PartialShape{1});
auto key_past = std::make_shared<ov::intel_gpu::op::ReadValue>(key_variable);
auto value_past = std::make_shared<ov::intel_gpu::op::ReadValue>(key_variable);
auto axis = std::make_shared<ov::op::v0::Constant>(ov::element::i64, ov::Shape{}, 0);
auto key_gather_past = std::make_shared<ov::op::v8::Gather>(key_past, beam_idx, axis);
auto value_gather_past = std::make_shared<ov::op::v8::Gather>(value_past, beam_idx, axis);
auto key_cache = std::make_shared<ov::intel_gpu::op::KVCache>(key_gather_past, key_token_param, key_variable, 2, ov::element::f32);
auto value_cache = std::make_shared<ov::intel_gpu::op::KVCache>(value_gather_past, value_token_param, value_variable, 2, ov::element::f32);
auto axes = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{1}, axes_val);
auto key_unsqueeze = std::make_shared<ov::op::v0::Unsqueeze>(key_cache, axes);
auto value_unsqueeze = std::make_shared<ov::op::v0::Unsqueeze>(value_cache, axes);
auto target_shape = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{5}, target_shape_kv);
auto key_broadcast = std::make_shared<ov::op::v3::Broadcast>(key_unsqueeze, target_shape, ov::op::BroadcastType::BIDIRECTIONAL);
auto value_broadcast = std::make_shared<ov::op::v3::Broadcast>(value_unsqueeze, target_shape, ov::op::BroadcastType::BIDIRECTIONAL);
auto pattern = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{4}, pattern_shape);
auto key_reshape = std::make_shared<ov::op::v1::Reshape>(key_broadcast, pattern, true);
auto value_reshape = std::make_shared<ov::op::v1::Reshape>(value_broadcast, pattern, true);
auto inputs = ov::OutputVector{input_q, key_reshape, value_reshape};
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(inputs, is_causal, in0_order, in1_order, in2_order, out_order);
model = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_q, key_token_param, value_token_param, beam_idx });
manager.register_pass<UnsqueezeBroadcastReshapeSDPAFusion>();
}
{
auto input_q = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto key_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{-1, -1, 2, 32}, ov::element::f32, "v0"});
auto value_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{-1, -1, 2, 32}, ov::element::f32, "v1"});
auto key_token_param = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{-1, -1, 2, 32});
auto value_token_param = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{-1, -1, 2, 32});
auto beam_idx = std::make_shared<ov::op::v0::Parameter>(ov::element::i32, ov::PartialShape{1});
auto key_past = std::make_shared<ov::intel_gpu::op::ReadValue>(key_variable);
auto value_past = std::make_shared<ov::intel_gpu::op::ReadValue>(key_variable);
auto axis = std::make_shared<ov::op::v0::Constant>(ov::element::i64, ov::Shape{}, 0);
auto key_gather_past = std::make_shared<ov::op::v8::Gather>(key_past, beam_idx, axis);
auto value_gather_past = std::make_shared<ov::op::v8::Gather>(value_past, beam_idx, axis);
auto key_cache = std::make_shared<ov::intel_gpu::op::KVCache>(key_gather_past, key_token_param, key_variable, 2, ov::element::f32);
auto value_cache = std::make_shared<ov::intel_gpu::op::KVCache>(value_gather_past, value_token_param, value_variable, 2, ov::element::f32);
auto inputs = ov::OutputVector{input_q, key_cache, value_cache};
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(inputs, is_causal, in0_order, in1_order, in2_order, out_order);
model_ref = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_q, key_token_param, value_token_param, beam_idx });
comparator.enable(FunctionsComparator::ATTRIBUTES);
}
}
TEST_F(TransformationTestsF, UnsqueezeBroadReshapeSDPAFusion2) {
std::vector<int64_t> in0_order = {0, 2, 1, 3};
std::vector<int64_t> in1_order = {0, 2, 1, 3};
std::vector<int64_t> in2_order = {0, 2, 1, 3};
std::vector<int64_t> out_order = {0, 1, 2, 3};
std::vector<int64_t> axes_val = {2};
std::vector<int64_t> pattern_shape = {0, 32, -1, 32};
const bool is_causal = true;
{
auto input_q = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto key_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{-1, 8, -1, 32}, ov::element::f32, "v0"});
auto value_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{-1, 8, -1, 32}, ov::element::f32, "v1"});
auto key_token_param = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{-1, 8, -1, 32});
auto value_token_param = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{-1, 8, -1, 32});
auto beam_idx = std::make_shared<ov::op::v0::Parameter>(ov::element::i32, ov::PartialShape{1});
auto key_past = std::make_shared<ov::intel_gpu::op::ReadValue>(key_variable);
auto value_past = std::make_shared<ov::intel_gpu::op::ReadValue>(key_variable);
auto axis = std::make_shared<ov::op::v0::Constant>(ov::element::i64, ov::Shape{}, 0);
auto key_gather_past = std::make_shared<ov::op::v8::Gather>(key_past, beam_idx, axis);
auto value_gather_past = std::make_shared<ov::op::v8::Gather>(value_past, beam_idx, axis);
auto key_cache = std::make_shared<ov::intel_gpu::op::KVCache>(key_gather_past, key_token_param, key_variable, 2, ov::element::f32);
auto value_cache = std::make_shared<ov::intel_gpu::op::KVCache>(value_gather_past, value_token_param, value_variable, 2, ov::element::f32);
auto axes = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{1}, axes_val);
auto key_unsqueeze = std::make_shared<ov::op::v0::Unsqueeze>(key_cache, axes);
auto value_unsqueeze = std::make_shared<ov::op::v0::Unsqueeze>(value_cache, axes);
auto abs_param = std::make_shared<ov::op::v0::Parameter>(ov::element::i32, ov::Shape{5});
auto abs = std::make_shared<ov::op::v0::Abs>(abs_param);
auto key_broadcast = std::make_shared<ov::op::v3::Broadcast>(key_unsqueeze, abs, ov::op::BroadcastType::BIDIRECTIONAL);
auto value_broadcast = std::make_shared<ov::op::v3::Broadcast>(value_unsqueeze, abs, ov::op::BroadcastType::BIDIRECTIONAL);
auto pattern = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{4}, pattern_shape);
auto key_reshape = std::make_shared<ov::op::v1::Reshape>(key_broadcast, pattern, true);
auto value_reshape = std::make_shared<ov::op::v1::Reshape>(value_broadcast, pattern, true);
auto inputs = ov::OutputVector{input_q, key_cache, value_cache};
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(inputs, is_causal, in0_order, in1_order, in2_order, out_order);
model = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_q, key_token_param, value_token_param, beam_idx });
manager.register_pass<UnsqueezeBroadcastReshapeSDPAFusion>();
}
{
auto input_q = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto key_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{-1, 8, -1, 32}, ov::element::f32, "v0"});
auto value_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{-1, 8, -1, 32}, ov::element::f32, "v1"});
auto key_token_param = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{-1, 8, -1, 32});
auto value_token_param = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{-1, 8, -1, 32});
auto beam_idx = std::make_shared<ov::op::v0::Parameter>(ov::element::i32, ov::PartialShape{1});
auto key_past = std::make_shared<ov::intel_gpu::op::ReadValue>(key_variable);
auto value_past = std::make_shared<ov::intel_gpu::op::ReadValue>(key_variable);
auto axis = std::make_shared<ov::op::v0::Constant>(ov::element::i64, ov::Shape{}, 0);
auto key_gather_past = std::make_shared<ov::op::v8::Gather>(key_past, beam_idx, axis);
auto value_gather_past = std::make_shared<ov::op::v8::Gather>(value_past, beam_idx, axis);
auto key_cache = std::make_shared<ov::intel_gpu::op::KVCache>(key_gather_past, key_token_param, key_variable, 2, ov::element::f32);
auto value_cache = std::make_shared<ov::intel_gpu::op::KVCache>(value_gather_past, value_token_param, value_variable, 2, ov::element::f32);
auto inputs = ov::OutputVector{input_q, key_cache, value_cache};
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(inputs, is_causal, in0_order, in1_order, in2_order, out_order);
model_ref = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_q, key_token_param, value_token_param, beam_idx });
comparator.enable(FunctionsComparator::ATTRIBUTES);
}
}
TEST_F(TransformationTestsF, UnsqueezeBroadReshapeSDPAFusion3) {
std::vector<int64_t> in0_order = {0, 2, 1, 3};
std::vector<int64_t> in1_order = {0, 2, 1, 3};
std::vector<int64_t> in2_order = {0, 2, 1, 3};
std::vector<int64_t> out_order = {0, 1, 2, 3};
std::vector<int64_t> axes_val = {-2};
std::vector<int32_t> target_shape_kv = {1, 1, 1, 16, 1};
std::vector<int64_t> pattern_shape = {0, 0, -1, 32};
const bool is_causal = true;
{
auto input_q = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto input_k = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{-1, -1, -1, 32});
auto input_v = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{-1, -1, -1, 32});
auto unsqueeze_const = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{1}, axes_val);
auto key_unsqueeze = std::make_shared<ov::op::v0::Unsqueeze>(input_k, unsqueeze_const);
auto value_unsqueeze = std::make_shared<ov::op::v0::Unsqueeze>(input_v, unsqueeze_const);
auto broadcast_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{5}, target_shape_kv);
auto key_broadcast = std::make_shared<ov::op::v3::Broadcast>(key_unsqueeze, broadcast_const, ov::op::BroadcastType::BIDIRECTIONAL);
auto value_broadcast = std::make_shared<ov::op::v3::Broadcast>(value_unsqueeze, broadcast_const, ov::op::BroadcastType::BIDIRECTIONAL);
auto reshape_const = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{4}, pattern_shape);
auto key_reshape = std::make_shared<ov::op::v1::Reshape>(key_broadcast, reshape_const, true);
auto value_reshape = std::make_shared<ov::op::v1::Reshape>(value_broadcast, reshape_const, true);
auto inputs = ov::OutputVector{input_q, key_reshape, value_reshape};
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(inputs, is_causal, in0_order, in1_order, in2_order, out_order);
model = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_q, input_k, input_v });
manager.register_pass<UnsqueezeBroadcastReshapeSDPAFusion>();
}
{
model_ref = model->clone();
comparator.enable(FunctionsComparator::ATTRIBUTES);
}
}
TEST_F(TransformationTestsF, UnsqueezeBroadReshapeSDPAFusion4) {
std::vector<int64_t> in0_order = {0, 2, 1, 3};
std::vector<int64_t> in1_order = {0, 2, 1, 3};
std::vector<int64_t> in2_order = {0, 2, 1, 3};
std::vector<int64_t> out_order = {0, 1, 2, 3};
std::vector<int64_t> axes_val = {-2};
std::vector<int32_t> target_shape_k = {1, 1, 1, 16, 1};
std::vector<int32_t> target_shape_v = {1, 1, 1, 32, 1};
std::vector<int64_t> pattern_shape = {0, 0, 32, 32};
const bool is_causal = true;
{
auto input_q = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(4));
auto key_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{-1, -1, 2, 32}, ov::element::f32, "v0"});
auto value_variable = std::make_shared<ov::op::util::Variable>(ov::op::util::VariableInfo{{-1, -1, 2, 32}, ov::element::f32, "v1"});
auto key_token_param = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{-1, -1, 2, 32});
auto value_token_param = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape{-1, -1, 2, 32});
auto beam_idx = std::make_shared<ov::op::v0::Parameter>(ov::element::i32, ov::PartialShape{1});
auto key_past = std::make_shared<ov::intel_gpu::op::ReadValue>(key_variable);
auto value_past = std::make_shared<ov::intel_gpu::op::ReadValue>(key_variable);
auto axis = std::make_shared<ov::op::v0::Constant>(ov::element::i64, ov::Shape{}, 0);
auto key_gather_past = std::make_shared<ov::op::v8::Gather>(key_past, beam_idx, axis);
auto value_gather_past = std::make_shared<ov::op::v8::Gather>(value_past, beam_idx, axis);
auto key_cache = std::make_shared<ov::intel_gpu::op::KVCache>(key_gather_past, key_token_param, key_variable, 2, ov::element::f32);
auto value_cache = std::make_shared<ov::intel_gpu::op::KVCache>(value_gather_past, value_token_param, value_variable, 2, ov::element::f32);
auto axes = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{1}, axes_val);
auto key_unsqueeze = std::make_shared<ov::op::v0::Unsqueeze>(key_cache, axes);
auto value_unsqueeze = std::make_shared<ov::op::v0::Unsqueeze>(value_cache, axes);
auto target_shape_key = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{5}, target_shape_k);
auto target_shape_value = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{5}, target_shape_v);
auto key_broadcast = std::make_shared<ov::op::v3::Broadcast>(key_unsqueeze, target_shape_key, ov::op::BroadcastType::BIDIRECTIONAL);
auto value_broadcast = std::make_shared<ov::op::v3::Broadcast>(value_unsqueeze, target_shape_value, ov::op::BroadcastType::BIDIRECTIONAL);
auto pattern = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{4}, pattern_shape);
auto key_reshape = std::make_shared<ov::op::v1::Reshape>(key_broadcast, pattern, true);
auto value_reshape = std::make_shared<ov::op::v1::Reshape>(value_broadcast, pattern, true);
auto inputs = ov::OutputVector{input_q, key_reshape, value_reshape};
auto sdpa = std::make_shared<ov::intel_gpu::op::SDPA>(inputs, is_causal, in0_order, in1_order, in2_order, out_order);
model = std::make_shared<ov::Model>(ov::NodeVector{ sdpa }, ov::ParameterVector{ input_q, key_token_param, value_token_param, beam_idx });
manager.register_pass<UnsqueezeBroadcastReshapeSDPAFusion>();
}
{
model_ref = model->clone();
comparator.enable(FunctionsComparator::ATTRIBUTES);
}
}
} // namespace intel_gpu
} // namespace test
} // namespace ov

2
src/plugins/intel_gpu/thirdparty/onednn_gpu vendored

@ -1 +1 @@
Subproject commit bab7bf42914073e443276e9efb0ddc5c93d6f52c
Subproject commit a7ed0bcbdefd0accc2c59d3030b7f8b21af12f6b

9
src/plugins/template/tests/functional/op_reference/squeeze.cpp
View File

@ -116,6 +116,15 @@ std::vector<SqueezeParams> generateParamsForSqueeze() {
using T2 = typename element_type_traits<Axes_T>::value_type;
std::vector<SqueezeParams> params{
SqueezeParams(Shape{1, 4, 1, 1, 2},
Shape{4, 1, 2},
IO_T,
IO_T,
std::vector<T1>{1, 2, 3, 4, 5, 6, 7, 8},
std::vector<T1>{1, 2, 3, 4, 5, 6, 7, 8},
Shape{4},
Axes_T,
std::vector<T2>{0, 1, 3, -1}),
SqueezeParams(Shape{1, 4, 1, 1, 2},
Shape{4, 1, 2},
IO_T,

56
src/tests/functional/plugin/shared/include/subgraph_tests/rotary_pos_emb.hpp Normal file
View File

@ -0,0 +1,56 @@
// Copyright (C) 2018-2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#pragma once
#include "shared_test_classes/subgraph/rotary_pos_emb.hpp"
namespace ov {
namespace test {
inline void CheckNumberOfNodesWithType(std::shared_ptr<const ov::Model> function,
const std::unordered_set<std::string>& nodeTypes,
size_t expectedCount) {
ASSERT_NE(nullptr, function);
int num_ops = 0;
for (const auto& node : function->get_ordered_ops()) {
const auto& rt_info = node->get_rt_info();
const auto layer_type = rt_info.find("layerType")->second.as<std::string>();
if (nodeTypes.count(layer_type)) {
num_ops++;
}
}
ASSERT_EQ(num_ops, expectedCount);
}
TEST_P(RoPETestLlama2, CompareWithRefs) {
SKIP_IF_CURRENT_TEST_IS_DISABLED();
run();
auto function = compiledModel.get_runtime_model();
CheckNumberOfNodesWithType(function, {"RoPE"}, 1);
};
TEST_P(RoPETestChatGLM, CompareWithRefs) {
SKIP_IF_CURRENT_TEST_IS_DISABLED();
run();
auto function = compiledModel.get_runtime_model();
CheckNumberOfNodesWithType(function, {"RoPE"}, 1);
};
TEST_P(RoPETestQwen7b, CompareWithRefs) {
SKIP_IF_CURRENT_TEST_IS_DISABLED();
run();
auto function = compiledModel.get_runtime_model();
CheckNumberOfNodesWithType(function, {"RoPE"}, 1);
};
TEST_P(RoPETestGPTJ, CompareWithRefs) {
SKIP_IF_CURRENT_TEST_IS_DISABLED();
run();
auto function = compiledModel.get_runtime_model();
CheckNumberOfNodesWithType(function, {"RoPE"}, 1);
};
} // namespace test
} // namespace ov

67
src/tests/functional/shared_test_classes/include/shared_test_classes/subgraph/rotary_pos_emb.hpp Normal file
View File

@ -0,0 +1,67 @@
// 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 {
class RoPETestLlama2 : public SubgraphBaseTest, public testing::WithParamInterface<std::string> {
private:
ov::OutputVector makeCosSinCache(int max_position_embeddings, int rotary_ndims);
std::shared_ptr<ov::Model> buildROPE_Llama2(int batch,
int seq_length,
int max_position_embeddings,
int num_head,
int ndims);
ov::Tensor create_i32_tensor(const ov::Shape& shape, int start, int step = 1);
protected:
void generate_inputs(const std::vector<ov::Shape>& targetInputStaticShapes) override;
void SetUp() override;
public:
static std::string getTestCaseName(const testing::TestParamInfo<std::string>& obj);
};
class RoPETestChatGLM : public SubgraphBaseTest, public testing::WithParamInterface<std::string> {
private:
std::shared_ptr<ov::Model> buildROPE_ChatGLM(int batch, int head_cnt, int rotary_dims);
ov::Tensor create_i32_tensor(const ov::Shape& shape, int start, int step = 1);
protected:
void generate_inputs(const std::vector<ov::Shape>& targetInputStaticShapes) override;
void SetUp() override;
public:
static std::string getTestCaseName(const testing::TestParamInfo<std::string>& obj);
};
class RoPETestQwen7b : public SubgraphBaseTest, public testing::WithParamInterface<std::tuple<bool, std::string>> {
private:
std::shared_ptr<ov::Model> buildROPE_QWen7b(bool specialReshape);
protected:
void generate_inputs(const std::vector<ov::Shape>& targetInputStaticShapes) override;
void SetUp() override;
public:
static std::string getTestCaseName(const testing::TestParamInfo<std::tuple<bool, std::string>>& obj);
};
class RoPETestGPTJ : public SubgraphBaseTest, public testing::WithParamInterface<std::tuple<bool, std::string>> {
private:
std::shared_ptr<ov::Model> buildROPE_GPTJ(int num_head,
int hidden_dims,
int rotary_dims,
bool hasShapeOf);
protected:
void generate_inputs(const std::vector<ov::Shape>& targetInputStaticShapes) override;
void SetUp() override;
public:
static std::string getTestCaseName(const testing::TestParamInfo<std::tuple<bool, std::string>>& obj);
};
} // namespace test
} // namespace ov

4
src/tests/functional/shared_test_classes/src/base/utils/ranges.cpp
View File

@ -85,12 +85,12 @@ void ModelRange::find_mode_ranges(const std::shared_ptr<ov::Model>& model) {
std::cout << ex.what() << std::endl;
#endif
}
// #ifndef NDEBUG
#ifndef NDEBUG
std::cout << "RANGE FOR PARAMETER: " << param->get_friendly_name()
<< " start from: " << std::to_string(data->start_from) << " range: " << std::to_string(data->range)
<< " resolution: " << std::to_string(data->resolution) << " seed: " << std::to_string(data->seed)
<< std::endl;
// #endif
#endif
std::string range_id = get_range_id(param);
node_ranges[range_id] = data;

590
src/tests/functional/shared_test_classes/src/subgraph/rotary_pos_emb.cpp Normal file
View File

@ -0,0 +1,590 @@
// Copyright (C) 2018-2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "common_test_utils/ov_tensor_utils.hpp"
#include "shared_test_classes/subgraph/rotary_pos_emb.hpp"
#include "transformations/utils/gen_pattern.hpp"
using namespace ov::gen_pattern;
using namespace ov;
namespace ov {
namespace test {
ov::OutputVector RoPETestLlama2::makeCosSinCache(int max_position_embeddings, int rotary_ndims) {
std::vector<float> lut_sin(max_position_embeddings * rotary_ndims, 0.0f);
std::vector<float> lut_cos(max_position_embeddings * rotary_ndims, 0.0f);
// rotate_half style cos/sin table:
// y1 = cos(m*xita_i) * x1 - sin(m*xita_i) * x2
// y2 = cos(m*xita_i) * x2 + sin(m*xita_i) * x1
//
for (int i = 0, k = 0; i < rotary_ndims; i += 2, k++) {
auto xita_i = 1.0 / std::pow(10000.0, static_cast<double>(i) / rotary_ndims);
float* psin = lut_sin.data();
float* pcos = lut_cos.data();
for (int m = 0; m < max_position_embeddings; m++, psin += rotary_ndims, pcos += rotary_ndims) {
auto vsin = std::sin(xita_i * m);
auto vcos = std::cos(xita_i * m);
pcos[k] = pcos[k + rotary_ndims / 2] = vcos;
psin[k] = psin[k + rotary_ndims / 2] = vsin;
}
}
auto shape = ov::Shape({1, 1, static_cast<size_t>(max_position_embeddings), static_cast<size_t>(rotary_ndims)});
auto Cos = makeConst(ov::element::f32, shape, lut_cos);
auto Sin = makeConst(ov::element::f32, shape, lut_sin);
return {Cos, Sin};
}
std::shared_ptr<ov::Model> RoPETestLlama2::buildROPE_Llama2(int batch,
int seq_length,
int max_position_embeddings,
int num_head,
int ndims) {
auto input = std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{batch, -1, num_head, ndims});
auto pos_id_end = std::make_shared<ov::opset1::Parameter>(ov::element::i32, ov::Shape{});
auto pos_ids = std::make_shared<ov::opset1::Parameter>(ov::element::i32, PartialShape{1, -1});
auto cos_sin_cache = makeCosSinCache(max_position_embeddings, ndims);
auto Constant582 = cos_sin_cache[0];
auto Constant585 = cos_sin_cache[1];
// concat KV length
auto transpose_Transpose = makeOP<ov::op::v1::Transpose>({input, {0, 2, 1, 3}});
auto slice_Unsqueeze_426 = makeOP<ov::op::v0::Unsqueeze>({pos_id_end, 0});
auto ScatterUpdate_152236 = makeOP<ov::op::v3::ScatterUpdate>({{0, 0, 0}, {2}, slice_Unsqueeze_426, {0}});
auto slice_Slice = makeOP<ov::op::v1::StridedSlice>({Constant582, {0, 0, 0}, ScatterUpdate_152236, {1, 1, 1}},
{{"begin_mask", {1, 1, 0}},
{"end_mask", {1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto squeeze_Squeeze = makeOP<ov::op::v0::Squeeze>({slice_Slice, 1});
auto squeeze_Squeeze_435 = makeOP<ov::op::v0::Squeeze>({squeeze_Squeeze, 0});
auto index_441_Gather = makeOP<ov::op::v8::Gather>({squeeze_Squeeze_435, pos_ids, 0}, {{"batch_dims", 0}});
auto unsqueeze_Unsqueeze = makeOP<ov::op::v0::Unsqueeze>({index_441_Gather, 1});
auto mul_Multiply =
makeOP<ov::op::v1::Multiply>({transpose_Transpose, unsqueeze_Unsqueeze}, {{"auto_broadcast", "numpy"}});
auto size_ShapeOf_448 = makeOP<ov::op::v3::ShapeOf>({transpose_Transpose}, {{"output_type", "i32"}});
auto size_Gather_450 = makeOP<ov::op::v8::Gather>({size_ShapeOf_448, 3, 0}, {{"batch_dims", 0}});
auto floor_divide_Divide =
makeOP<ov::op::v1::Divide>({size_Gather_450, 2}, {{"auto_broadcast", "numpy"}, {"m_pythondiv", true}});
auto floor_divide_Floor = makeOP<ov::op::v0::Floor>({floor_divide_Divide});
auto slice_Unsqueeze_452 = makeOP<ov::op::v0::Unsqueeze>({floor_divide_Floor, 0});
auto ScatterUpdate_152312 = makeOP<ov::op::v3::ScatterUpdate>({{0, 0, 0, 0}, {3}, slice_Unsqueeze_452, {0}});
auto slice_Slice_459 = makeOP<ov::op::v1::StridedSlice>(
{transpose_Transpose, ScatterUpdate_152312, {0ll, 0ll, 0ll, LLONG_MAX}, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto Constant_182988 = makeConst(element::f32,
ov::Shape({
1,
1,
1,
1,
}),
{-1.000000f});
auto neg_Multiply = makeOP<ov::op::v1::Multiply>({slice_Slice_459, Constant_182988}, {{"auto_broadcast", "numpy"}});
auto ScatterUpdate_152368 = makeOP<ov::op::v3::ScatterUpdate>({{0, 0, 0, 0}, {3}, slice_Unsqueeze_452, {0}});
auto slice_Slice2 =
makeOP<ov::op::v1::StridedSlice>({transpose_Transpose, {0, 0, 0, 0}, ScatterUpdate_152368, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto cat_Concat = makeOP<ov::op::v0::Concat>({neg_Multiply, slice_Slice2}, {{"axis", -1}});
auto ScatterUpdate_152421 = makeOP<ov::op::v3::ScatterUpdate>({{0, 0, 0}, {2}, slice_Unsqueeze_426, {0}});
auto slice_Slice_433 = makeOP<ov::op::v1::StridedSlice>({Constant585, {0, 0, 0}, ScatterUpdate_152421, {1, 1, 1}},
{{"begin_mask", {1, 1, 0}},
{"end_mask", {1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto squeeze_Squeeze_436 = makeOP<ov::op::v0::Squeeze>({slice_Slice_433, 1});
auto squeeze_Squeeze_437 = makeOP<ov::op::v0::Squeeze>({squeeze_Squeeze_436, 0});
auto index_446_Gather = makeOP<ov::op::v8::Gather>({squeeze_Squeeze_437, pos_ids, 0}, {{"batch_dims", 0}});
auto unsqueeze_Unsqueeze_447 = makeOP<ov::op::v0::Unsqueeze>({index_446_Gather, 1});
auto mul_Multiply_463 =
makeOP<ov::op::v1::Multiply>({cat_Concat, unsqueeze_Unsqueeze_447}, {{"auto_broadcast", "numpy"}});
auto add_Add = makeOP<ov::op::v1::Add>({mul_Multiply, mul_Multiply_463}, {{"auto_broadcast", "numpy"}});
return std::make_shared<ov::Model>(ov::NodeVector{add_Add}, ov::ParameterVector{input, pos_id_end, pos_ids});
}
ov::Tensor RoPETestLlama2::create_i32_tensor(const ov::Shape& shape, int start, int step) {
auto tensor = ov::Tensor(ov::element::i32, shape);
auto* ptr = static_cast<int32_t*>(tensor.data());
for (size_t i = 0; i < tensor.get_size(); i++) {
ptr[i] = start;
start += step;
}
return tensor;
}
void RoPETestLlama2::generate_inputs(const std::vector<ov::Shape>& targetInputStaticShapes) {
const auto& funcInputs = function->inputs();
const int position_id_start = 15;
auto& input_shape = targetInputStaticShapes[0];
auto seq_length = input_shape[1];
ov::test::utils::InputGenerateData in_data;
in_data.start_from = -1;
in_data.range = 2;
in_data.resolution = 32768;
ov::Tensor t_input = utils::create_and_fill_tensor(funcInputs[0].get_element_type(), input_shape, in_data);
ov::Tensor t_position_id_end = create_i32_tensor(ov::Shape({}), position_id_start + seq_length);
ov::Tensor t_position_ids = create_i32_tensor(ov::Shape({1, seq_length}), position_id_start);
inputs.clear();
inputs.insert({funcInputs[0].get_node_shared_ptr(), t_input});
inputs.insert({funcInputs[1].get_node_shared_ptr(), t_position_id_end});
inputs.insert({funcInputs[2].get_node_shared_ptr(), t_position_ids});
}
void RoPETestLlama2::SetUp() {
targetDevice = this->GetParam();
const int batch = 2;
const int seq_length = 7;
const size_t max_position_embeddings = 2048;
const size_t ndims = 128;
const size_t num_head = 32;
InputShape inpShape = {{batch, seq_length, num_head, ndims}, {{batch, seq_length, num_head, ndims}}};
init_input_shapes({inpShape});
function = buildROPE_Llama2(batch, seq_length, max_position_embeddings, num_head, ndims);
}
std::string RoPETestLlama2::getTestCaseName(const testing::TestParamInfo<std::string>& obj) {
std::string targetDevice = obj.param;
std::ostringstream result;
result << "targetDevice=" << targetDevice;
return result.str();
}
std::shared_ptr<ov::Model> RoPETestChatGLM::buildROPE_ChatGLM(int batch, int head_cnt, int rotary_dims) {
auto input =
std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{-1, batch, 4096 + 256 + 256});
auto cos_sin_cache = std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{32768, 32, 2});
auto position_ids = std::make_shared<ov::opset1::Parameter>(ov::element::i32, PartialShape{-1, -1});
auto __module_transformer_index_67_Gather =
makeOP<opset8::Gather>({cos_sin_cache, position_ids, 0}, {{"batch_dims", 0}});
auto __module_transformer_transpose_Transpose =
makeOP<opset1::Transpose>({__module_transformer_index_67_Gather, {1, 0, 2, 3}});
auto size_ShapeOf_110 =
makeOP<opset3::ShapeOf>({__module_transformer_transpose_Transpose}, {{"output_type", "i32"}});
auto __getitem___Gather = makeOP<opset8::Gather>({size_ShapeOf_110, -2, 0}, {{"batch_dims", 0}});
auto mul_Multiply = makeOP<opset1::Multiply>({__getitem___Gather, 2}, {{"auto_broadcast", "numpy"}});
auto slice_Unsqueeze_112 = makeOP<opset1::Unsqueeze>({mul_Multiply, 0});
auto floordiv_Divide =
makeOP<opset1::Divide>({mul_Multiply, 2}, {{"auto_broadcast", "numpy"}, {"m_pythondiv", true}});
auto floordiv_Floor = makeOP<opset1::Floor>({floordiv_Divide});
auto ListConstruct_126_Reshape_2 = makeOP<opset1::Reshape>({floordiv_Floor, {-1}}, {{"special_zero", false}});
auto ListUnpack_321 = makeOP<opset1::VariadicSplit>({input, -1, {4096, 256, 256}});
auto view_Reshape =
makeOP<opset1::Reshape>({ListUnpack_321->output(0), {0, 0, 32, 128}}, {{"special_zero", true}});
auto ScatterUpdate_229053 = makeOP<opset3::ScatterUpdate>({{0, 0, 0, 0}, {3}, slice_Unsqueeze_112, {0}});
auto slice_Slice_357 =
makeOP<opset1::StridedSlice>({view_Reshape, {0, 0, 0, 0}, ScatterUpdate_229053, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto size_ShapeOf_346 = makeOP<opset3::ShapeOf>({view_Reshape}, {{"output_type", "i32"}});
auto size_Gather_348 = makeOP<opset8::Gather>({size_ShapeOf_346, 0, 0}, {{"batch_dims", 0}});
auto ListConstruct_372_Reshape = makeOP<opset1::Reshape>({size_Gather_348, {-1}}, {{"special_zero", false}});
auto size_Gather_351 = makeOP<opset8::Gather>({size_ShapeOf_346, {2}, 0}, {{"batch_dims", 0}});
auto ListConstruct_372_Concat =
makeOP<opset1::Concat>({ListConstruct_372_Reshape, {-1}, size_Gather_351, ListConstruct_126_Reshape_2, {2}},
{{"axis", 0}});
auto reshape_Reshape_373 =
makeOP<opset1::Reshape>({slice_Slice_357, ListConstruct_372_Concat}, {{"special_zero", false}});
auto select_Gather_381 = makeOP<opset8::Gather>({reshape_Reshape_373, 0, -1}, {{"batch_dims", 0}});
auto slice_Unsqueeze_367 = makeOP<opset1::Unsqueeze>({size_Gather_348, 0});
auto slice_Slice_369 =
makeOP<opset1::StridedSlice>({__module_transformer_transpose_Transpose, {0}, slice_Unsqueeze_367, {1}},
{{"begin_mask", {0}},
{"end_mask", {0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto size_ShapeOf_374 = makeOP<opset3::ShapeOf>({reshape_Reshape_373}, {{"output_type", "i32"}});
auto size_Gather_376 = makeOP<opset8::Gather>({size_ShapeOf_374, {3}, 0}, {{"batch_dims", 0}});
auto ListConstruct_379_Concat =
makeOP<opset1::Concat>({ListConstruct_372_Reshape, {-1}, {1}, size_Gather_376, {2}}, {{"axis", 0}});
auto view_Reshape_380 =
makeOP<opset1::Reshape>({slice_Slice_369, ListConstruct_379_Concat}, {{"special_zero", false}});
auto select_Gather_382 = makeOP<opset8::Gather>({view_Reshape_380, 0, -1}, {{"batch_dims", 0}});
auto mul_Multiply_383 =
makeOP<opset1::Multiply>({select_Gather_381, select_Gather_382}, {{"auto_broadcast", "numpy"}});
auto select_Gather_384 = makeOP<opset8::Gather>({reshape_Reshape_373, 1, -1}, {{"batch_dims", 0}});
auto select_Gather_385 = makeOP<opset8::Gather>({view_Reshape_380, 1, -1}, {{"batch_dims", 0}});
auto mul_Multiply_386 =
makeOP<opset1::Multiply>({select_Gather_384, select_Gather_385}, {{"auto_broadcast", "numpy"}});
auto sub_Subtract_389 =
makeOP<opset1::Subtract>({mul_Multiply_383, mul_Multiply_386}, {{"auto_broadcast", "numpy"}});
auto Unsqueeze_62716 = makeOP<opset1::Unsqueeze>({sub_Subtract_389, -1});
auto mul_Multiply_391 =
makeOP<opset1::Multiply>({select_Gather_384, select_Gather_382}, {{"auto_broadcast", "numpy"}});
auto mul_Multiply_393 =
makeOP<opset1::Multiply>({select_Gather_381, select_Gather_385}, {{"auto_broadcast", "numpy"}});
auto add_Add_396 = makeOP<opset1::Add>({mul_Multiply_391, mul_Multiply_393}, {{"auto_broadcast", "numpy"}});
auto Unsqueeze_62717 = makeOP<opset1::Unsqueeze>({add_Add_396, -1});
auto stack_401 = makeOP<opset1::Concat>({Unsqueeze_62716, Unsqueeze_62717}, {{"axis", -1}});
auto flatten_ShapeOf_402 = makeOP<opset3::ShapeOf>({stack_401}, {{"output_type", "i32"}});
auto flatten_Slice_417 = makeOP<opset1::StridedSlice>({flatten_ShapeOf_402, {0}, {3}, {1}},
{{"begin_mask", {0}},
{"end_mask", {0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto flatten_Concat_420 = makeOP<opset1::Concat>({flatten_Slice_417, {-1}}, {{"axis", 0}});
auto flatten_Reshape_421 = makeOP<opset1::Reshape>({stack_401, flatten_Concat_420}, {{"special_zero", true}});
auto ScatterUpdate_229067 = makeOP<opset3::ScatterUpdate>({{0, 0, 0, 0}, {3}, slice_Unsqueeze_112, {0}});
auto slice_Slice_363 =
makeOP<opset1::StridedSlice>({view_Reshape, ScatterUpdate_229067, {0, 0, 0, INT_MAX}, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto cat_Concat_425 = makeOP<opset1::Concat>({flatten_Reshape_421, slice_Slice_363}, {{"axis", -1}});
return std::make_shared<ov::Model>(ov::NodeVector{cat_Concat_425},
ov::ParameterVector{input, cos_sin_cache, position_ids});
}
ov::Tensor RoPETestChatGLM::create_i32_tensor(const ov::Shape& shape, int start, int step) {
auto tensor = ov::Tensor(ov::element::i32, shape);
auto* ptr = static_cast<int32_t*>(tensor.data());
for (size_t i = 0; i < tensor.get_size(); i++) {
ptr[i] = start;
start += step;
}
return tensor;
}
void RoPETestChatGLM::generate_inputs(const std::vector<ov::Shape>& targetInputStaticShapes) {
const auto& funcInputs = function->inputs();
auto& input_shape = targetInputStaticShapes[0];
auto seq_length = input_shape[0];
ov::Tensor t_input =
utils::create_and_fill_tensor(funcInputs[0].get_element_type(), input_shape, 2, -1.0f, 32768);
ov::Tensor t_cos_sin_cache =
utils::create_and_fill_tensor(funcInputs[1].get_element_type(), {32768, 32, 2}, 2, -1.0f, 32768);
ov::Tensor t_position_ids = create_i32_tensor(ov::Shape({1, seq_length}), 15);
inputs.clear();
inputs.insert({funcInputs[0].get_node_shared_ptr(), t_input});
inputs.insert({funcInputs[1].get_node_shared_ptr(), t_cos_sin_cache});
inputs.insert({funcInputs[2].get_node_shared_ptr(), t_position_ids});
}
void RoPETestChatGLM::SetUp() {
targetDevice = this->GetParam();
const int batch = 2;
const int seq_length = 7;
const int num_head = 32;
const int rotary_dims = 64;
InputShape inpShape = {{-1, batch, 4096 + 256 + 256}, {{seq_length, batch, 4096 + 256 + 256}}};
init_input_shapes({inpShape});
function = buildROPE_ChatGLM(batch, num_head, rotary_dims);
}
std::string RoPETestChatGLM::getTestCaseName(const testing::TestParamInfo<std::string>& obj) {
std::string targetDevice = obj.param;
std::ostringstream result;
result << "targetDevice=" << targetDevice;
return result.str();
}
std::shared_ptr<ov::Model> RoPETestQwen7b::buildROPE_QWen7b(bool specialReshape) {
auto input =
std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{-1, -1, 4096 + 4096 + 4096});
auto cos_cache = std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{1, -1, 1, 128});
auto sin_cache = std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{1, -1, 1, 128});
auto ListUnpack_389_VariadicSplit = makeOP<opset1::VariadicSplit>({input, 2, {4096, 4096, -1}});
auto view_Reshape = makeOP<opset1::Reshape>({ListUnpack_389_VariadicSplit->output(0), {0, 0, 32, 128}},
{{"special_zero", true}});
auto size_ShapeOf_414 = makeOP<opset3::ShapeOf>({view_Reshape}, {{"output_type", "i32"}});
auto size_Gather_416 = makeOP<opset8::Gather>({size_ShapeOf_414, 1, 0}, {{"batch_dims", 0}});
auto neg_Multiply = makeOP<opset1::Multiply>({size_Gather_416, -1}, {{"auto_broadcast", "numpy"}});
auto slice_Unsqueeze_422 = makeOP<opset1::Unsqueeze>({neg_Multiply, 0});
auto ScatterUpdate_261437 = makeOP<opset3::ScatterUpdate>({{0, 0}, {1}, slice_Unsqueeze_422, {0}});
auto slice_Slice_425 = makeOP<opset1::StridedSlice>({cos_cache, ScatterUpdate_261437, {0ll, LLONG_MAX}, {1, 1}},
{{"begin_mask", {1, 0}},
{"end_mask", {1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto slice_Slice_431 =
makeOP<opset1::StridedSlice>({slice_Slice_425, {0, 0, 0}, {0ll, 0ll, LLONG_MAX}, {1, 1, 1}},
{{"begin_mask", {1, 1, 0}},
{"end_mask", {1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto slice_Slice_437 =
makeOP<opset1::StridedSlice>({slice_Slice_431, {0, 0, 0, 0}, {0ll, 0ll, 0ll, LLONG_MAX}, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto size_ShapeOf_462 = makeOP<opset3::ShapeOf>({slice_Slice_437}, {{"output_type", "i32"}});
auto size_Gather_464 = makeOP<opset8::Gather>({size_ShapeOf_462, {3}, 0}, {{"batch_dims", 0}});
auto ScatterUpdate_261533 = makeOP<opset3::ScatterUpdate>({{0, 0, 0, 0}, {3}, size_Gather_464, {0}});
auto slice_Slice_470 =
makeOP<opset1::StridedSlice>({view_Reshape, {0, 0, 0, 0}, ScatterUpdate_261533, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto mul_Multiply = makeOP<opset1::Multiply>({slice_Slice_470, slice_Slice_437}, {{"auto_broadcast", "numpy"}});
auto size_ShapeOf_478 = makeOP<opset3::ShapeOf>({slice_Slice_470}, {{"output_type", "i32"}});
auto Gather_239390 = makeOP<opset8::Gather>({size_ShapeOf_478, {0, 1, 2}, 0}, {{"batch_dims", 0}});
auto size_Gather_489 = makeOP<opset8::Gather>({size_ShapeOf_478, 3, 0}, {{"batch_dims", 0}});
auto floor_divide_Divide =
makeOP<opset1::Divide>({size_Gather_489, 2}, {{"auto_broadcast", "numpy"}, {"m_pythondiv", true}});
auto floor_divide_Floor = makeOP<opset1::Floor>({floor_divide_Divide});
auto ListConstruct_493_Reshape_3 =
makeOP<opset1::Reshape>({floor_divide_Floor, {-1}}, {{"special_zero", false}});
auto ListConstruct_493_Concat =
makeOP<opset1::Concat>({Gather_239390, {2}, ListConstruct_493_Reshape_3}, {{"axis", 0}});
std::shared_ptr<ov::Node> reshape_Reshape = nullptr;
if (specialReshape) {
reshape_Reshape = makeOP<opset1::Reshape>({slice_Slice_470, {0, 0, 32, 2, 64}}, {{"special_zero", true}});
} else {
reshape_Reshape =
makeOP<opset1::Reshape>({slice_Slice_470, ListConstruct_493_Concat}, {{"special_zero", false}});
}
auto ListUnpack_496_Split = makeOP<opset1::Split>({reshape_Reshape, -2}, {{"num_splits", 2}});
auto ListUnpack_496_Squeeze_0 = makeOP<opset1::Squeeze>({ListUnpack_496_Split->output(1), -2});
auto Constant_296840_compressed = makeConst(element::f16,
ov::Shape({
1,
1,
1,
1,
}),
{-1});
auto Constant_296840 = makeOP<opset1::Convert>({Constant_296840_compressed}, {{"destination_type", "f32"}});
auto neg_Multiply_499 =
makeOP<opset1::Multiply>({ListUnpack_496_Squeeze_0, Constant_296840}, {{"auto_broadcast", "numpy"}});
auto ListUnpack_496_Squeeze = makeOP<opset1::Squeeze>({ListUnpack_496_Split->output(0), -2});
auto cat_Concat = makeOP<opset1::Concat>({neg_Multiply_499, ListUnpack_496_Squeeze}, {{"axis", -1}});
auto slice_Slice_449 = makeOP<opset1::StridedSlice>({sin_cache, ScatterUpdate_261437, {0ll, LLONG_MAX}, {1, 1}},
{{"begin_mask", {1, 0}},
{"end_mask", {1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto slice_Slice_455 =
makeOP<opset1::StridedSlice>({slice_Slice_449, {0, 0, 0}, {0ll, 0ll, LLONG_MAX}, {1, 1, 1}},
{{"begin_mask", {1, 1, 0}},
{"end_mask", {1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto slice_Slice_461 =
makeOP<opset1::StridedSlice>({slice_Slice_455, {0, 0, 0, 0}, {0ll, 0ll, 0ll, LLONG_MAX}, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto mul_Multiply_503 = makeOP<opset1::Multiply>({cat_Concat, slice_Slice_461}, {{"auto_broadcast", "numpy"}});
auto add_Add = makeOP<opset1::Add>({mul_Multiply, mul_Multiply_503}, {{"auto_broadcast", "numpy"}});
return std::make_shared<ov::Model>(ov::NodeVector{add_Add}, ov::ParameterVector{input, cos_cache, sin_cache});
}
void RoPETestQwen7b::generate_inputs(const std::vector<ov::Shape>& targetInputStaticShapes) {
const auto& funcInputs = function->inputs();
auto& input_shape = targetInputStaticShapes[0];
ov::Tensor t_input =
utils::create_and_fill_tensor(funcInputs[0].get_element_type(), input_shape, 2, -1.0f, 32768);
ov::Tensor t_cos_cache =
utils::create_and_fill_tensor(funcInputs[1].get_element_type(), {1, 4096, 1, 128}, 2, -1.0f, 32768);
ov::Tensor t_sin_cache =
utils::create_and_fill_tensor(funcInputs[1].get_element_type(), {1, 4096, 1, 128}, 2, -1.0f, 32768);
inputs.clear();
inputs.insert({funcInputs[0].get_node_shared_ptr(), t_input});
inputs.insert({funcInputs[1].get_node_shared_ptr(), t_cos_cache});
inputs.insert({funcInputs[2].get_node_shared_ptr(), t_sin_cache});
}
void RoPETestQwen7b::SetUp() {
bool specialReshape;
std::tie(specialReshape, targetDevice) = this->GetParam();
const int batch = 2;
const int seq_length = 7;
InputShape inpShape = {{batch, -1, 4096 + 4096 + 4096}, {{batch, seq_length, 4096 + 4096 + 4096}}};
init_input_shapes({inpShape});
function = buildROPE_QWen7b(specialReshape);
}
std::string RoPETestQwen7b::getTestCaseName(const testing::TestParamInfo<std::tuple<bool, std::string>>& obj) {
bool specialReshape;
std::string targetDevice;
std::tie(specialReshape, targetDevice) = obj.param;
std::ostringstream result;
result << "specialReshape=" << specialReshape << "_"
<< "targetDevice=" << targetDevice;
return result.str();
}
std::shared_ptr<ov::Model> RoPETestGPTJ::buildROPE_GPTJ(int num_head,
int hidden_dims,
int rotary_dims,
bool hasShapeOf) {
auto int32_max = std::numeric_limits<std::int32_t>::max();
auto input =
std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{-1, -1, num_head, hidden_dims});
auto sincos = std::make_shared<ov::opset1::Parameter>(ov::element::f32, PartialShape{-1, -1, rotary_dims});
auto slice_Slice_965 =
makeOP<ov::op::v1::StridedSlice>({input, {0, 0, 0, 0}, {0, 0, 0, rotary_dims}, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
slice_Slice_965->set_friendly_name("slice_Slice_965");
auto varsplit = makeOP<ov::op::v1::VariadicSplit>({sincos, -1, {rotary_dims / 2, -1}});
varsplit->set_output_size(2);
varsplit->set_friendly_name("varsplit");
auto unsqueeze_sin = makeOP<opset1::Unsqueeze>({varsplit->output(0), 2});
auto unsqueeze_cos = makeOP<opset1::Unsqueeze>({varsplit->output(1), 2});
std::vector<int32_t> gather_idx(rotary_dims, 1);
int32_t v = 0;
for (size_t i = 0; i < gather_idx.size(); i += 2, v++) {
gather_idx[i] = v;
gather_idx[i + 1] = v;
}
auto const_idx = makeConst(ov::element::i32, ov::Shape({static_cast<size_t>(rotary_dims)}), gather_idx);
auto constant_155588 = makeConst(element::f32,
ov::Shape({
1,
1,
1,
1,
}),
{-1.000000f});
auto repeat_interleave_sin = makeOP<opset8::Gather>({unsqueeze_sin, const_idx, 3}, {{"batch_dims", 0}});
auto repeat_interleave_cos = makeOP<opset8::Gather>({unsqueeze_cos, const_idx, 3}, {{"batch_dims", 0}});
repeat_interleave_sin->set_friendly_name("repeat_interleave_sin");
repeat_interleave_cos->set_friendly_name("repeat_interleave_cos");
// x interleave (-x[:,:,:, 1::2], x[:,:,:, 0::2])
auto slice_Slice_1174 =
makeOP<ov::op::v1::StridedSlice>({slice_Slice_965, {0, 0, 0, 1}, {0, 0, 0, int32_max}, {1, 1, 1, 2}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto neg_Multiply_1177 =
makeOP<opset1::Multiply>({slice_Slice_1174, constant_155588}, {{"auto_broadcast", "numpy"}});
auto Unsqueeze_65524 = makeOP<opset1::Unsqueeze>({neg_Multiply_1177, -1});
auto slice_Slice_1168 =
makeOP<ov::op::v1::StridedSlice>({slice_Slice_965, {0, 0, 0, 0}, {0, 0, 0, int32_max}, {1, 1, 1, 2}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto Unsqueeze_65525 = makeOP<opset1::Unsqueeze>({slice_Slice_1168, -1});
auto stack_1182 = makeOP<opset1::Concat>({Unsqueeze_65524, Unsqueeze_65525}, {{"axis", -1}});
auto flatten_Reshape_1198 =
makeOP<opset1::Reshape>({stack_1182, {0, 0, num_head, rotary_dims}}, {{"special_zero", true}});
// x*cos [B,L,H,ndims]
auto mul_cos =
makeOP<opset1::Multiply>({slice_Slice_965, repeat_interleave_cos}, {{"auto_broadcast", "numpy"}});
mul_cos->set_friendly_name("mul_cos");
auto mul_sin =
makeOP<opset1::Multiply>({flatten_Reshape_1198, repeat_interleave_sin}, {{"auto_broadcast", "numpy"}});
// *cos + *sin
auto rotary_emb = makeOP<opset1::Add>({mul_cos, mul_sin}, {{"auto_broadcast", "numpy"}});
auto slice_Slice_971 =
makeOP<ov::op::v1::StridedSlice>({input, {0, 0, 0, rotary_dims}, {0, 0, 0, int32_max}, {1, 1, 1, 1}},
{{"begin_mask", {1, 1, 1, 0}},
{"end_mask", {1, 1, 1, 0}},
{"new_axis_mask", {}},
{"shrink_axis_mask", {}},
{"ellipsis_mask", {}}});
auto cat_Concat_1211 = makeOP<opset1::Concat>({rotary_emb, slice_Slice_971}, {{"axis", -1}});
auto permute_Transpose_1213 = makeOP<opset1::Transpose>({cat_Concat_1211, {0, 2, 1, 3}});
ov::NodeVector model_output = {permute_Transpose_1213};
if (hasShapeOf) {
auto shapeOf = makeOP<opset1::ShapeOf>({rotary_emb}, {{"output_type", "i32"}});
auto gather = makeOP<opset8::Gather>({shapeOf, {1}, 0}, {{"batch_dims", 0}});
model_output.push_back(gather);
}
return std::make_shared<ov::Model>(model_output, ov::ParameterVector{input, sincos});
}
void RoPETestGPTJ::generate_inputs(const std::vector<ov::Shape>& targetInputStaticShapes) {
const auto& funcInputs = function->inputs();
auto& input_shape = targetInputStaticShapes[0];
auto& sincos_shape = targetInputStaticShapes[1];
ov::Tensor t_input =
utils::create_and_fill_tensor(funcInputs[0].get_element_type(), input_shape, 2, -1.0f, 32768);
ov::Tensor t_cos_sin_cache =
utils::create_and_fill_tensor(funcInputs[1].get_element_type(), sincos_shape, 2, -1.0f, 32768);
inputs.clear();
inputs.insert({funcInputs[0].get_node_shared_ptr(), t_input});
inputs.insert({funcInputs[1].get_node_shared_ptr(), t_cos_sin_cache});
}
std::string RoPETestGPTJ::getTestCaseName(const testing::TestParamInfo<std::tuple<bool, std::string>>& obj) {
bool hasShapeOf;
std::string targetDevice;
std::tie(hasShapeOf, targetDevice) = obj.param;
std::ostringstream result;
result << "hasShapeOf=" << hasShapeOf << "_"
<< "targetDevice=" << targetDevice;
return result.str();
}
void RoPETestGPTJ::SetUp() {
bool hasShapeOf;
std::tie(hasShapeOf, targetDevice) = this->GetParam();
const int batch = 2;
const int seq_length = 7;
const int num_head = 16;
const int hidden_dims = 256;
const int rotary_dims = 64;
InputShape input = {{batch, seq_length, num_head, hidden_dims}, {{batch, seq_length, num_head, hidden_dims}}};
InputShape sincos = {{batch, seq_length, rotary_dims}, {{batch, seq_length, rotary_dims}}};
init_input_shapes({input, sincos});
function = buildROPE_GPTJ(num_head, hidden_dims, rotary_dims, hasShapeOf);
}
} // namespace test
} // namespace ov

1
tests/layer_tests/pytorch_tests/test_squeeze.py
View File

@ -39,7 +39,6 @@ class TestSqueeze(PytorchLayerTest):
pytest.xfail(reason="export fails if dim is not provided")
self._test(*self.create_model(dim), ie_device, precision, ir_version, dynamic_shapes=dynamic_shapes)
@pytest.mark.xfail(reason='OpenVINO squeeze does not support dimension is not equal to 1.')
@pytest.mark.parametrize("dim", [-1, 2])
@pytest.mark.nightly
@pytest.mark.precommit

6
tests/layer_tests/tensorflow_tests/test_tf_Rint.py
View File

@ -1,6 +1,8 @@
# Copyright (C) 2018-2024 Intel Corporation
# SPDX-License-Identifier: Apache-2.0
import platform
import numpy as np
import pytest
import tensorflow as tf
@ -32,6 +34,10 @@ class TestRint(CommonTFLayerTest):
@pytest.mark.parametrize("input_type", [np.float32, np.float64])
@pytest.mark.precommit
@pytest.mark.nightly
@pytest.mark.xfail(condition=platform.system() in ('Darwin', 'Linux') and platform.machine() in ['arm', 'armv7l',
'aarch64',
'arm64', 'ARM64'],
reason='Ticket - 126314, 132699')
def test_rint_basic(self, input_shape, input_type, ie_device, precision,
ir_version, temp_dir, use_legacy_frontend):
self._test(*self.create_tf_rint_net(input_shape, input_type),