[GPU] eye-like operation (#11937)

* eye-like operation

single layer tests
unit tests

* [GPU] adjust eye implementation to the current GPU operation design

src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp

@@ -241,6 +241,7 @@ REGISTER_FACTORY(v9, SoftSign)
 REGISTER_FACTORY(v9, ROIAlign);
 REGISTER_FACTORY(v9, RDFT);
 REGISTER_FACTORY(v9, IRDFT);
+REGISTER_FACTORY(v9, Eye);
 
 // --------------------------- Supported internal ops --------------------------- //
 REGISTER_FACTORY(internal, NonMaxSuppressionIEInternal);

src/plugins/intel_gpu/include/intel_gpu/primitives/eye.hpp

@@ -0,0 +1,42 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#pragma once
+#include "primitive.hpp"
+
+namespace cldnn {
+/// @addtogroup cpp_api C++ API
+/// @{
+/// @addtogroup cpp_topology Network Topology
+/// @{
+/// @addtogroup cpp_primitives Primitives
+/// @{
+
+/// @brief
+/// @details Construct identity matrix or batch fo them
+struct eye : public primitive_base<eye> {
+    CLDNN_DECLARE_PRIMITIVE(eye)
+
+    /// @brief Constructs eye primitive.
+    /// @param id This primitive id.
+    /// @param inputs List of primitive ids.
+    /// @param output_shape Tensor output shape
+    /// @param ext_prim_id Primitive extra id (friendly name)
+    /// @param shift Eye diagonal
+    /// @param output_type Tensor output type
+    eye(const primitive_id& id,
+        const std::vector<primitive_id>& inputs,
+        const tensor& output_shape,
+        const int32_t shift,
+        const cldnn::data_types output_type)
+        : primitive_base{id, inputs, padding(), optional_data_type(output_type)},
+          output_shape{output_shape},
+          shift{shift} {}
+
+    tensor output_shape;
+    int32_t shift;
+};
+/// @}
+/// @}
+/// @}
+}  // namespace cldnn

src/plugins/intel_gpu/src/graph/eye.cpp

@@ -0,0 +1,40 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include <eye_inst.h>
+#include <json_object.h>
+
+#include <sstream>
+
+#include "primitive_type_base.h"
+
+namespace cldnn {
+
+primitive_type_id eye::type_id() {
+    static primitive_type_base<eye> instance;
+    return &instance;
+}
+
+eye_inst::typed_primitive_inst(network& network, eye_node const& node) : parent(network, node) {}
+
+layout eye_inst::calc_output_layout(eye_node const& node, const kernel_impl_params&) {
+    auto primitive = node.get_primitive();
+    return {*(primitive->output_data_type), node.input().get_output_layout().format, primitive->output_shape};
+}
+
+std::string eye_inst::to_string(eye_node const& node) {
+    auto node_info = node.desc_to_json();
+    json_composite eye_info;
+    eye_info.add("rows id", node.get_dependency(0).id());
+    eye_info.add("cols id", node.get_dependency(1).id());
+    eye_info.add("diagInd id", node.get_dependency(2).id());
+    if (node.get_dependencies().size() == 4)
+        eye_info.add("batchShape id", node.get_dependency(3).id());
+    node_info->add("slice info", eye_info);
+    std::stringstream primitive_description;
+    node_info->dump(primitive_description);
+    return primitive_description.str();
+}
+
+}  // namespace cldnn

src/plugins/intel_gpu/src/graph/impls/ocl/eye.cpp

@@ -0,0 +1,79 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include <data_inst.h>
+#include <eye/eye_kernel_ref.h>
+#include <eye/eye_kernel_selector.h>
+#include <eye_inst.h>
+
+#include <algorithm>
+#include <cstddef>
+#include <impls/implementation_map.hpp>
+#include <intel_gpu/runtime/error_handler.hpp>
+#include <vector>
+
+#include "primitive_base.hpp"
+
+namespace cldnn {
+namespace ocl {
+
+struct eye_impl : typed_primitive_impl_ocl<eye> {
+    using parent = typed_primitive_impl_ocl<eye>;
+    using parent::parent;
+
+    std::unique_ptr<primitive_impl> clone() const override {
+        return make_unique<eye_impl>(*this);
+    }
+
+    static primitive_impl* create(const eye_node& arg, const kernel_impl_params& impl_param) {
+        auto params = get_default_params<kernel_selector::eye_params>(impl_param);
+        auto op_params = get_default_optional_params<kernel_selector::eye_optional_params>(arg.get_program());
+
+        auto primitive = arg.get_primitive();
+        params.diagonal_index = primitive->shift;
+
+        auto& kernel_selector = kernel_selector::eye_kernel_selector::Instance();
+        auto best_kernels = kernel_selector.GetBestKernels(params, op_params);
+
+        CLDNN_ERROR_BOOL(arg.id(),
+                         "Best_kernel.empty()",
+                         best_kernels.empty(),
+                         "Cannot find a proper kernel with this arguments");
+
+        return new eye_impl(arg, best_kernels[0]);
+    }
+};
+
+namespace detail {
+
+attach_eye_impl::attach_eye_impl() {
+    const std::vector<data_types> types{data_types::f16,
+                                        data_types::f32,
+                                        data_types::i8,
+                                        data_types::u8,
+                                        data_types::i32,
+                                        data_types::i64};
+    const std::vector<format::type> formats{
+        format::bfyx,
+        format::b_fs_yx_fsv16,
+        format::b_fs_yx_fsv32,
+        format::bs_fs_yx_bsv16_fsv16,
+        format::bs_fs_yx_bsv32_fsv32,
+        format::bs_fs_yx_bsv32_fsv16,
+
+        format::bfzyx,
+        format::b_fs_zyx_fsv16,
+        format::b_fs_zyx_fsv32,
+        format::bs_fs_zyx_bsv16_fsv32,
+        format::bs_fs_zyx_bsv16_fsv16,
+        format::bs_fs_zyx_bsv32_fsv32,
+        format::bs_fs_zyx_bsv32_fsv16,
+    };
+    implementation_map<eye>::add(impl_types::ocl, eye_impl::create, types, formats);
+}
+
+}  // namespace detail
+
+}  // namespace ocl
+}  // namespace cldnn

src/plugins/intel_gpu/src/graph/impls/ocl/register.cpp

@@ -92,6 +92,7 @@ void register_implementations() {
     REGISTER_OCL(convert_color);
     REGISTER_OCL(count_nonzero);
     REGISTER_OCL(gather_nonzero);
+    REGISTER_OCL(eye);
 }
 
 }  // namespace ocl

src/plugins/intel_gpu/src/graph/impls/ocl/register.hpp

@@ -27,6 +27,7 @@
 #include "intel_gpu/primitives/experimental_detectron_prior_grid_generator.hpp"
 #include "intel_gpu/primitives/experimental_detectron_roi_feature_extractor.hpp"
 #include "intel_gpu/primitives/experimental_detectron_topk_rois.hpp"
+#include "intel_gpu/primitives/eye.hpp"
 #include "intel_gpu/primitives/fully_connected.hpp"
 #include "intel_gpu/primitives/gather.hpp"
 #include "intel_gpu/primitives/gather_elements.hpp"
@@ -73,6 +74,7 @@
 #include "intel_gpu/primitives/strided_slice.hpp"
 #include "intel_gpu/primitives/tile.hpp"
 #include "intel_gpu/primitives/non_zero.hpp"
+#include "intel_gpu/primitives/eye.hpp"
 
 namespace cldnn {
 namespace ocl {
@@ -168,6 +170,7 @@ REGISTER_OCL(extract_image_patches);
 REGISTER_OCL(convert_color);
 REGISTER_OCL(count_nonzero);
 REGISTER_OCL(gather_nonzero);
+REGISTER_OCL(eye);
 
 #undef REGISTER_OCL
 

src/plugins/intel_gpu/src/graph/include/eye_inst.h

@@ -0,0 +1,41 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma once
+#include <intel_gpu/primitives/eye.hpp>
+#include <intel_gpu/runtime/error_handler.hpp>
+
+#include "primitive_inst.h"
+
+namespace cldnn {
+
+template <>
+struct typed_program_node<eye> : public typed_program_node_base<eye> {
+    using parent = typed_program_node_base<eye>;
+
+public:
+    using parent::parent;
+
+    program_node& input(std::size_t index = 0) const {
+        return get_dependency(index);
+    }
+};
+
+using eye_node = typed_program_node<eye>;
+
+template <>
+class typed_primitive_inst<eye> : public typed_primitive_inst_base<eye> {
+    using parent = typed_primitive_inst_base<eye>;
+
+public:
+    static layout calc_output_layout(eye_node const& node, const kernel_impl_params& impl_param);
+    static std::string to_string(eye_node const& node);
+
+public:
+    typed_primitive_inst(network& network, eye_node const& desc);
+};
+
+using eye_inst = typed_primitive_inst<eye>;
+
+}  // namespace cldnn

src/plugins/intel_gpu/src/graph/program.cpp

@@ -1431,7 +1431,8 @@ void program::set_layout_optimizer_attributes(layout_optimizer& lo) {
             prim.type() != cldnn::bucketize::type_id() &&
             prim.type() != cldnn::roll::type_id() &&
             prim.type() != cldnn::prior_box::type_id() &&
-            prim.type() != cldnn::resample::type_id()) {
+            prim.type() != cldnn::resample::type_id() &&
+            prim.type() != cldnn::eye::type_id()) {
             can_use_fsv16 = false;
         }
 
@@ -1466,7 +1467,8 @@ void program::set_layout_optimizer_attributes(layout_optimizer& lo) {
             prim.type() != cldnn::bucketize::type_id() &&
             prim.type() != cldnn::roll::type_id() &&
             prim.type() != cldnn::resample::type_id() &&
-            prim.type() != cldnn::prior_box::type_id()) {
+            prim.type() != cldnn::prior_box::type_id() &&
+            prim.type() != cldnn::eye::type_id()) {
             can_use_bs_fs_yx_bsv16_fsv16 = false;
         }
     }

src/plugins/intel_gpu/src/kernel_selector/common/common_types.h

@@ -90,6 +90,7 @@ enum class KernelType {
     ADAPTIVE_POOLING,
     REVERSE,
     PRIOR_BOX,
+    EYE,
 };
 
 ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/eye/eye_kernel_ref.cpp

@@ -0,0 +1,84 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include "eye_kernel_ref.h"
+
+#include <kernel_selector_utils.h>
+
+#include <vector>
+
+namespace kernel_selector {
+
+KernelsData EyeKernelRef::GetKernelsData(const Params& params, const optional_params& options) const {
+    if (!Validate(params, options)) {
+        return {};
+    }
+    KernelData kernel_data = KernelData::Default<eye_params>(params);
+    const eye_params& new_params = dynamic_cast<eye_params&>(*kernel_data.params.get());
+    auto dispatch_data = SetDefault(new_params, options);
+    auto entry_point = GetEntryPoint(kernelName, new_params.layerID, params, options);
+    auto slice_specific_jit = GetJitConstants(new_params);
+    auto jit = CreateJit(kernelName, slice_specific_jit, entry_point);
+
+    FillCLKernelData(kernel_data.kernels[0], dispatch_data, params.engineInfo, kernelName, jit, entry_point);
+    return {kernel_data};
+}
+
+KernelsPriority EyeKernelRef::GetKernelsPriority(const Params& /*params*/, const optional_params& /*options*/) const {
+    return DONT_USE_IF_HAVE_SOMETHING_ELSE;
+}
+
+ParamsKey EyeKernelRef::GetSupportedKey() const {
+    ParamsKey k;
+    k.EnableInputDataType(Datatype::INT32);
+    k.EnableInputDataType(Datatype::INT64);
+    // Inputs for Eye operation can be only integer types but in case of blocked layout and when output is float,
+    // the inputs are transformed to f32.
+    // That is the reason why the operation below is present
+    k.EnableInputDataType(Datatype::F32);
+    k.EnableOutputDataType(Datatype::F16);
+    k.EnableOutputDataType(Datatype::F32);
+    k.EnableOutputDataType(Datatype::INT32);
+    k.EnableOutputDataType(Datatype::INT8);
+    k.EnableOutputDataType(Datatype::UINT8);
+    k.EnableOutputDataType(Datatype::INT64);
+    k.EnableAllInputLayout();
+    k.EnableAllOutputLayout();
+    k.EnableTensorOffset();
+    k.EnableTensorPitches();
+    k.EnableBatching();
+    k.EnableDifferentTypes();
+    return k;
+}
+
+bool EyeKernelRef::Validate(const Params& p, const optional_params& o) const {
+    if (p.GetType() != KernelType::EYE || o.GetType() != KernelType::EYE) {
+        return false;
+    }
+
+    const eye_params& params = dynamic_cast<const eye_params&>(p);
+    if (params.inputs.empty())
+        return false;
+
+    return true;
+}
+
+JitConstants EyeKernelRef::GetJitConstants(const eye_params& params) const {
+    JitConstants jit = MakeBaseParamsJitConstants(params);
+    jit.AddConstant(MakeJitConstant("DIAGONAL", params.diagonal_index));
+    return jit;
+}
+
+CommonDispatchData EyeKernelRef::SetDefault(const eye_params& params, const optional_params&) const {
+    CommonDispatchData dispatchData;
+    dispatchData.gws = {params.outputs[0].Batch().v,
+                        params.outputs[0].Feature().v,
+                        params.outputs[0].Z().v * params.outputs[0].Y().v * params.outputs[0].X().v};
+
+    dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, params.engineInfo);
+
+    return dispatchData;
+}
+
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/eye/eye_kernel_ref.h

@@ -0,0 +1,36 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma once
+
+#include <vector>
+
+#include "kernel_base_opencl.h"
+
+namespace kernel_selector {
+
+struct eye_params : public base_params {
+    eye_params() : base_params(KernelType::EYE) {}
+
+    std::int32_t diagonal_index;
+};
+
+struct eye_optional_params : optional_params {
+    eye_optional_params() : optional_params(KernelType::EYE) {}
+};
+
+class EyeKernelRef : public KernelBaseOpenCL {
+public:
+    EyeKernelRef() : KernelBaseOpenCL{"eye_ref"} {}
+    KernelsData GetKernelsData(const Params& params, const optional_params& options) const override;
+    KernelsPriority GetKernelsPriority(const Params& params, const optional_params& options) const override;
+    ParamsKey GetSupportedKey() const override;
+    bool Validate(const Params& p, const optional_params& o) const override;
+
+private:
+    JitConstants GetJitConstants(const eye_params& params) const;
+    CommonDispatchData SetDefault(const eye_params& params, const optional_params&) const;
+};
+
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/eye/eye_kernel_selector.cpp

@@ -0,0 +1,18 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#include "eye_kernel_selector.h"
+
+#include "eye_kernel_ref.h"
+
+namespace kernel_selector {
+
+eye_kernel_selector::eye_kernel_selector() {
+    Attach<EyeKernelRef>();
+}
+
+KernelsData eye_kernel_selector::GetBestKernels(const Params& params, const optional_params& options) const {
+    return GetNaiveBestKernel(params, options, KernelType::EYE);
+}
+
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/eye/eye_kernel_selector.h

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

src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/eye_ref.cl

@@ -0,0 +1,29 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include "include/batch_headers/fetch_data.cl"
+
+KERNEL(eye_ref)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output) {
+    // Get the index of the current element to be processed
+    const uint b = get_global_id(0);
+    const uint f = get_global_id(1);
+#if OUTPUT_DIMS <= 4
+    const uint xy = get_global_id(2);
+    const uint y = xy / OUTPUT_SIZE_X;
+    const uint x = xy % OUTPUT_SIZE_X;
+    const uint idx = OUTPUT_GET_INDEX(b, f, y, x);
+#elif OUTPUT_DIMS == 5
+    const uint xyz = get_global_id(2);
+    const uint yx = xyz % (OUTPUT_SIZE_X * OUTPUT_SIZE_Y);
+    const uint z = xyz / (OUTPUT_SIZE_X * OUTPUT_SIZE_Y);
+    const uint y = yx / OUTPUT_SIZE_X;
+    const uint x = yx % OUTPUT_SIZE_X;
+    const uint idx = OUTPUT_GET_INDEX(b, f, z, y, x);
+#endif
+    if (x - DIAGONAL == y) {
+        output[idx] = 1;
+    } else {
+        output[idx] = 0;
+    }
+}

src/plugins/intel_gpu/src/plugin/ops/eye.cpp

@@ -0,0 +1,60 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include "ngraph/op/eye.hpp"
+
+#include <memory>
+
+#include "intel_gpu/plugin/common_utils.hpp"
+#include "intel_gpu/plugin/program.hpp"
+#include "intel_gpu/primitives/eye.hpp"
+#include "intel_gpu/runtime/layout.hpp"
+
+namespace ov {
+namespace intel_gpu {
+
+namespace {
+
+static void CreateEyeOp(Program& p, const std::shared_ptr<ngraph::op::v9::Eye>& op) {
+    validate_inputs_count(op, {3, 4});
+
+    const InferenceEngine::SizeVector& output_shapes = op->get_output_shape(0);
+    auto os_sz = output_shapes.size();
+    assert(2 <= os_sz && os_sz <= 5);
+
+    size_t dim_size = std::max(os_sz, static_cast<size_t>(4));
+    InferenceEngine::SizeVector dims(dim_size, 1);
+    for (size_t i = dim_size, j = os_sz; i > 0 && j > 0; --i, --j) {
+        dims[i - 1] = output_shapes[j - 1];
+    }
+    const ngraph::op::v0::Constant* constant = dynamic_cast<const ngraph::op::v0::Constant*>(op->get_input_node_ptr(2));
+    int32_t shift{};
+    switch (constant->get_element_type()) {
+    case ov::element::Type_t::i32:
+        shift = *constant->get_data_ptr<int32_t>();
+        break;
+    case ov::element::Type_t::i64:
+        shift = *constant->get_data_ptr<int64_t>();
+        break;
+    default:
+        throw std::runtime_error{"Input type can be only either i32 or i64"};
+        break;
+    }
+    auto input_primitives = p.GetInputPrimitiveIDs(op);
+    auto output_shape = tensor_from_dims(dims);
+    auto eye_prim = cldnn::eye(layer_type_name_ID(op),
+                               input_primitives,
+                               output_shape,
+                               shift,
+                               cldnn::element_type_to_data_type(op->get_out_type()));
+
+    p.add_primitive(*op, eye_prim);
+}
+
+}  // namespace
+
+REGISTER_FACTORY_IMPL(v9, Eye);
+
+}  // namespace intel_gpu
+}  // namespace ov

src/plugins/intel_gpu/src/plugin/transformations_pipeline.cpp

@@ -94,6 +94,7 @@
 #include <low_precision/multiply_to_group_convolution.hpp>
 #include <low_precision/strided_slice.hpp>
 #include <low_precision/network_helper.hpp>
+#include "transformations/op_conversions/eye_decomposition.hpp"
 
 #include "intel_gpu/plugin/itt.hpp"
 
@@ -132,6 +133,7 @@ void TransformationsPipeline::apply(std::shared_ptr<ov::Model> func) {
         manager.register_pass<ngraph::pass::InitNodeInfo>();
         manager.register_pass<EinsumDecomposition>();
         manager.register_pass<ngraph::pass::CommonOptimizations>();
+
         manager.register_pass<ngraph::pass::WrapInterpolateIntoTransposes>();
         manager.register_pass<ngraph::pass::TransposeSinking>();
 
@@ -188,7 +190,7 @@ void TransformationsPipeline::apply(std::shared_ptr<ov::Model> func) {
         manager.register_pass<ngraph::pass::ConvertPrecision>(convert_precision_list);
 
         auto pass_config = manager.get_pass_config();
-
+        pass_config->disable<ov::pass::EyeDecomposition>();
         pass_config->enable<ov::pass::ConvertCompressedOnlyToLegacy>();
 
         // SpaceToDepth/DepthToSpace node implementation supports only equal input/output tensors with rank <= 5

src/plugins/intel_gpu/tests/test_cases/eye.cpp

@@ -0,0 +1,207 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#include <algorithm>
+#include <intel_gpu/primitives/data.hpp>
+#include <intel_gpu/primitives/eye.hpp>
+#include <intel_gpu/primitives/input_layout.hpp>
+#include <random>
+#include <vector>
+
+#include "test_utils.h"
+
+using namespace cldnn;
+using namespace ::tests;
+
+namespace {
+
+template <class OutputType, class InputType>
+using eye_test_param = std::tuple<format,                    // Input and output format
+                                  InputType,                 // columns number
+                                  InputType,                 // rows number
+                                  InputType,                 // diagonal index
+                                  std::vector<InputType>,    // batch shape
+                                  std::vector<int32_t>,      // output shape
+                                  std::vector<OutputType>>;  // expected values
+
+template <class OutputType, class InputType>
+class EyeTest : public ::testing::TestWithParam<eye_test_param<OutputType, InputType>> {
+public:
+    void SetUp() override {
+        format fmt{format::bfyx};
+        InputType cols{};
+        InputType rows{};
+        InputType diag{};
+        std::vector<InputType> batch_shape;
+        std::vector<int32_t> output_shape;
+        std::vector<OutputType> expected_values;
+
+        std::tie(fmt, cols, rows, diag, batch_shape, output_shape, expected_values) = this->GetParam();
+
+        auto num_rows = engine_.allocate_memory({type_to_data_type<InputType>::value, fmt, tensor{1}});
+        set_values<InputType>(num_rows, {rows});
+        auto num_coloms = engine_.allocate_memory({type_to_data_type<InputType>::value, fmt, tensor{1}});
+        set_values<InputType>(num_coloms, {cols});
+        auto diagonal_index = engine_.allocate_memory({type_to_data_type<InputType>::value, fmt, tensor{1}});
+        set_values<InputType>(diagonal_index, {diag});
+
+        topology tp;
+        tp.add(data("num_rows", num_rows));
+        tp.add(data("num_columns", num_coloms));
+        tp.add(data("diagonal_index", diagonal_index));
+
+        auto batch_rank = batch_shape.size() == 3 ? 3 : 2;
+        auto oupput_fmt = batch_rank == 3 ? format::bfzyx : format::bfyx;
+        if (!batch_shape.empty()) {
+            auto batch = engine_.allocate_memory({type_to_data_type<InputType>::value, fmt, tensor{batch_rank}});
+            set_values<InputType>(batch, batch_shape);
+            tp.add(data("batch", batch));
+        }
+
+        std::string ouput_op_name;
+        if (fmt == format::bfyx || fmt == format::bfzyx) {
+            auto inputs = batch_shape.empty()
+                              ? std::vector<primitive_id>{"num_rows", "num_columns", "diagonal_index"}
+                              : std::vector<primitive_id>{"num_rows", "num_columns", "diagonal_index", "batch"};
+            ouput_op_name = "eye";
+            auto eye_primitive =
+                eye("eye", inputs, tensor{output_shape}, diag, type_to_data_type<OutputType>::value);
+            tp.add(std::move(eye_primitive));
+        } else {
+            tp.add(reorder("r_num_rows", "num_rows", fmt, type_to_data_type<InputType>::value));
+            tp.add(reorder("r_num_columns", "num_columns", fmt, type_to_data_type<InputType>::value));
+            tp.add(reorder("r_diagonal_index", "diagonal_index", fmt, type_to_data_type<InputType>::value));
+            if (!batch_shape.empty()) {
+                tp.add(reorder("r_batch", "batch", fmt, type_to_data_type<InputType>::value));
+            }
+            auto inputs = batch_shape.empty()
+                              ? std::vector<primitive_id>{"r_num_rows", "r_num_columns", "r_diagonal_index"}
+                              : std::vector<primitive_id>{"r_num_rows", "r_num_columns", "r_diagonal_index", "r_batch"};
+            auto eye_primitive =
+                eye("eye", inputs, tensor{output_shape}, diag, type_to_data_type<OutputType>::value);
+            tp.add(std::move(eye_primitive));
+            ouput_op_name = "output";
+            tp.add(reorder("output", "eye", oupput_fmt, type_to_data_type<OutputType>::value));
+        }
+
+        network network(engine_, tp);
+
+        auto outputs = network.execute();
+
+        EXPECT_EQ(outputs.size(), size_t(1));
+        EXPECT_EQ(outputs.begin()->first, ouput_op_name);
+
+        auto output = outputs.at(ouput_op_name).get_memory();
+
+        cldnn::mem_lock<OutputType> output_ptr(output, get_test_stream());
+
+        ASSERT_EQ(output_ptr.size(), expected_values.size());
+        for (size_t i = 0; i < output_ptr.size(); ++i)
+            EXPECT_TRUE(are_equal(expected_values[i], output_ptr[i], 2e-3));
+    }
+
+protected:
+    engine& engine_ = get_test_engine();
+};
+
+std::vector<format> four_d_formats{
+    format::bfyx,
+    format::b_fs_yx_fsv16,
+    format::b_fs_yx_fsv32,
+    format::bs_fs_yx_bsv16_fsv16,
+    format::bs_fs_yx_bsv32_fsv32,
+    format::bs_fs_yx_bsv32_fsv16,
+};
+
+using eye_test_4d_float_int32 = EyeTest<float, int32_t>;
+TEST_P(eye_test_4d_float_int32, eye_test_4d_float_int32) {}
+INSTANTIATE_TEST_SUITE_P(
+    eye_test_4d_float_int32,
+    eye_test_4d_float_int32,
+    testing::Combine(testing::ValuesIn(four_d_formats),
+                     testing::Values(2),
+                     testing::Values(3),
+                     testing::Values(0),
+                     testing::ValuesIn(std::vector<std::vector<int32_t>>{{}, {1}, {1, 1}, {1, 1, 1}}),
+                     testing::Values(std::vector<int32_t>{1, 1, 2, 3}),
+                     testing::Values(std::vector<float>{1, 0, 0, 1, 0, 0})));
+
+using eye_test_4d_int64_int32 = EyeTest<int64_t, int32_t>;
+TEST_P(eye_test_4d_int64_int32, eye_test_4d_int64_int32) {}
+INSTANTIATE_TEST_SUITE_P(
+    eye_test_4d_int64_int32,
+    eye_test_4d_int64_int32,
+    testing::Combine(testing::ValuesIn(four_d_formats),
+                     testing::Values(2),
+                     testing::Values(3),
+                     testing::Values(0),
+                     testing::ValuesIn(std::vector<std::vector<int32_t>>{{}, {1}, {1, 1}, {1, 1, 1}}),
+                     testing::Values(std::vector<int32_t>{1, 1, 2, 3}),
+                     testing::Values(std::vector<int64_t>{1, 0, 0, 1, 0, 0})));
+
+using eye_test_4d_u8_int64 = EyeTest<uint8_t, int64_t>;
+TEST_P(eye_test_4d_u8_int64, eye_test_4d_u8_int64) {}
+INSTANTIATE_TEST_SUITE_P(
+    eye_test_4d_u8_int64,
+    eye_test_4d_u8_int64,
+    testing::Combine(testing::ValuesIn(four_d_formats),
+                     testing::Values(4),
+                     testing::Values(3),
+                     testing::Values(-1),
+                     testing::ValuesIn(std::vector<std::vector<int64_t>>{{}, {1}, {1, 1}, {1, 1, 1}}),
+                     testing::Values(std::vector<int32_t>{1, 1, 4, 3}),
+                     testing::Values(std::vector<uint8_t>{0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0})));
+
+using eye_test_4d_i8_int64_no_diag = EyeTest<int8_t, int64_t>;
+TEST_P(eye_test_4d_i8_int64_no_diag, eye_test_4d_i8_int64_no_diag) {}
+INSTANTIATE_TEST_SUITE_P(
+    eye_test_4d_i8_int64_no_diag,
+    eye_test_4d_i8_int64_no_diag,
+    testing::Combine(testing::ValuesIn(four_d_formats),
+                     testing::Values(4),
+                     testing::Values(3),
+                     testing::Values(4),
+                     testing::ValuesIn(std::vector<std::vector<int64_t>>{{}, {1}, {1, 1}, {1, 1, 1}}),
+                     testing::Values(std::vector<int32_t>{1, 1, 4, 3}),
+                     testing::Values(std::vector<int8_t>{0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0})));
+
+using eye_test_4d_int32_int32_batch = EyeTest<int32_t, int32_t>;
+TEST_P(eye_test_4d_int32_int32_batch, eye_test_4d_int32_int32_batch) {}
+INSTANTIATE_TEST_SUITE_P(
+    eye_test_4d_int32_int32_batch,
+    eye_test_4d_int32_int32_batch,
+    testing::Combine(testing::ValuesIn(four_d_formats),
+                     testing::Values(2),
+                     testing::Values(2),
+                     testing::Values(1),
+                     testing::ValuesIn(std::vector<std::vector<int32_t>>{{2, 2}}),
+                     testing::Values(std::vector<int32_t>{2, 2, 2, 2}),
+                     testing::Values(std::vector<int32_t>{0, 1, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0})));
+
+std::vector<format> five_d_formats{
+    format::bfzyx,
+    format::b_fs_zyx_fsv16,
+    format::b_fs_zyx_fsv32,
+    format::bs_fs_zyx_bsv16_fsv32,
+    format::bs_fs_zyx_bsv16_fsv16,
+    format::bs_fs_zyx_bsv32_fsv32,
+    format::bs_fs_zyx_bsv32_fsv16,
+};
+
+using eye_test_5d_float_int32 = EyeTest<float, int32_t>;
+TEST_P(eye_test_5d_float_int32, eye_test_5d_float_int32) {}
+INSTANTIATE_TEST_SUITE_P(eye_test_5d_float_int32,
+                         eye_test_5d_float_int32,
+                         testing::Combine(testing::ValuesIn(five_d_formats),
+                                          testing::Values(2),
+                                          testing::Values(2),
+                                          testing::Values(0),
+                                          testing::ValuesIn(std::vector<std::vector<int32_t>>{{2, 2, 2}}),
+                                          testing::Values(std::vector<int32_t>{2, 2, 2, 2, 2}),
+                                          testing::Values(std::vector<float>{
+                                              1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1,
+
+                                              1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1,
+                                          })));
+
+}  // anonymous namespace

src/tests/functional/plugin/gpu/single_layer_tests/eye.cpp

@@ -0,0 +1,77 @@
+/// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include "shared_test_classes/single_layer/eye.hpp"
+
+using namespace LayerTestsDefinitions;
+
+namespace {
+
+TEST_P(EyeLayerTest, CompareWithRefs) {
+    SKIP_IF_CURRENT_TEST_IS_DISABLED()
+    Run();
+}
+
+const std::vector<ov::element::Type_t> netPrecisions =
+    {ElementType::f32, ElementType::f16, ElementType::i32, ElementType::i8, ElementType::u8, ElementType::i64};
+
+const std::vector<std::vector<int>> eyePars = {
+    // rows, cols, diag_shift
+    {3, 3, 0},
+    {3, 4, 1},
+    {4, 3, 1},
+    {3, 4, 0},
+    {4, 3, 0},
+    {3, 4, -1},
+    {4, 3, -1},
+    {3, 4, 10},
+    {4, 4, -2},
+};
+
+// dummy parameter to prevent empty set of test cases
+const std::vector<std::vector<int>> emptyBatchShape = {{0}};
+const std::vector<std::vector<int>> batchShapes1D = {{3}, {2}, {1}, {0}};
+const std::vector<std::vector<int>> batchShapes2D = {{3, 2}, {2, 1}, {0, 0}};
+const std::vector<std::vector<int>> batchShapes3D = {{3, 2, 1}, {1, 1, 1}};
+
+INSTANTIATE_TEST_SUITE_P(smoke_Eye2D_WithNonScalar_Test,
+                         EyeLayerTest,
+                         ::testing::Combine(::testing::ValuesIn(std::vector<std::vector<ov::Shape>>{{{1}, {1}, {1}}}),
+                                            ::testing::ValuesIn(emptyBatchShape),
+                                            ::testing::ValuesIn(eyePars),
+                                            ::testing::ValuesIn(netPrecisions),
+                                            ::testing::Values(CommonTestUtils::DEVICE_GPU)),
+                         EyeLayerTest::getTestCaseName);
+
+INSTANTIATE_TEST_SUITE_P(smoke_Eye_1DBatch_Test,
+                         EyeLayerTest,
+                         ::testing::Combine(::testing::ValuesIn(std::vector<std::vector<ov::Shape>>{
+                                                {{1}, {1}, {1}, {1}}}),
+                                            ::testing::ValuesIn(batchShapes1D),
+                                            ::testing::ValuesIn(eyePars),
+                                            ::testing::ValuesIn(netPrecisions),
+                                            ::testing::Values(CommonTestUtils::DEVICE_GPU)),
+                         EyeLayerTest::getTestCaseName);
+
+INSTANTIATE_TEST_SUITE_P(smoke_Eye_2DBatch_Test,
+                         EyeLayerTest,
+                         ::testing::Combine(::testing::ValuesIn(std::vector<std::vector<ov::Shape>>{
+                                                {{1}, {1}, {1}, {2}}}),
+                                            ::testing::ValuesIn(batchShapes2D),
+                                            ::testing::ValuesIn(eyePars),
+                                            ::testing::ValuesIn(netPrecisions),
+                                            ::testing::Values(CommonTestUtils::DEVICE_GPU)),
+                         EyeLayerTest::getTestCaseName);
+
+INSTANTIATE_TEST_SUITE_P(smoke_Eye_3DBatch_Test,
+                         EyeLayerTest,
+                         ::testing::Combine(::testing::ValuesIn(std::vector<std::vector<ov::Shape>>{
+                                                {{1}, {1}, {1}, {3}}}),
+                                            ::testing::ValuesIn(batchShapes3D),
+                                            ::testing::ValuesIn(eyePars),
+                                            ::testing::ValuesIn(netPrecisions),
+                                            ::testing::Values(CommonTestUtils::DEVICE_GPU)),
+                         EyeLayerTest::getTestCaseName);
+
+}  // namespace

src/tests/functional/plugin/shared/src/single_layer_tests/op_impl_check/single_op_graph.cpp

@@ -4,6 +4,7 @@
 
 #include <single_layer_tests/op_impl_check/op_impl_check.hpp>
 #include <single_layer_tests/op_impl_check/single_op_graph.hpp>
+#include <openvino/pass/constant_folding.hpp>
 
 namespace ov {
 namespace test {
@@ -368,14 +369,18 @@ std::shared_ptr<ov::Model> generate(const std::shared_ptr<ov::op::v3::ExtractIma
 }
 
 std::shared_ptr<ov::Model> generate(const std::shared_ptr<ov::op::v9::Eye> &node) {
-    const auto params = ngraph::builder::makeDynamicParams(ov::element::i64, {{1}, {1}, {1}, {3}});
-    const auto eye = std::make_shared<ov::op::v9::Eye>(params[0],
-                                                       params[1],
-                                                       params[2],
-                                                       params[3],
+    const auto rows = ngraph::builder::makeConstant<int64_t>(ov::element::i64, {1}, {3});
+    const auto cols = ngraph::builder::makeConstant<int64_t>(ov::element::i64, {1}, {4});
+    const auto diag = ngraph::builder::makeConstant<int64_t>(ov::element::i64, {1}, {0});
+    const auto batch = ngraph::builder::makeConstant<int64_t>(ov::element::i64, {3}, {2, 2, 2});
+    const auto eye = std::make_shared<ov::op::v9::Eye>(rows,
+                                                       cols,
+                                                       diag,
+                                                       batch,
                                                        ov::element::f32);
+    ov::pass::disable_constant_folding(eye);
     ov::ResultVector results{std::make_shared<ov::op::v0::Result>(eye)};
-    return std::make_shared<ov::Model>(results, params, "Eye");
+    return std::make_shared<ov::Model>(results, ngraph::ParameterVector{}, "Eye");
 }
 
 std::shared_ptr<ov::Model> generate(const std::shared_ptr<ov::op::v0::FakeQuantize> &node) {

src/tests/functional/shared_test_classes/include/shared_test_classes/single_layer/eye.hpp

@@ -0,0 +1,27 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#pragma once
+
+#include "shared_test_classes/base/layer_test_utils.hpp"
+
+namespace LayerTestsDefinitions {
+
+using ElementType = ov::element::Type_t;
+using TargetDevice = std::string;
+using LocalElementType = ov::element_type_traits<ngraph::element::i32>::value_type;
+
+using EyeLayerTestParams = std::tuple<std::vector<ov::Shape>,  // eye shape
+                                      std::vector<int>,        // output batch shape
+                                      std::vector<int>,        // eye params (rows, cols, diag_shift)
+                                      ElementType,             // Net precision
+                                      TargetDevice>;           // Device name
+
+class EyeLayerTest : public testing::WithParamInterface<EyeLayerTestParams>,
+                     virtual public LayerTestsUtils::LayerTestsCommon {
+public:
+    static std::string getTestCaseName(testing::TestParamInfo<EyeLayerTestParams> obj);
+    void SetUp() override;
+};
+
+}  // namespace LayerTestsDefinitions

src/tests/functional/shared_test_classes/src/single_layer/eye.cpp

@@ -0,0 +1,78 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#include "shared_test_classes/single_layer/eye.hpp"
+
+#include <common_test_utils/ov_tensor_utils.hpp>
+#include <ngraph/opsets/opset9.hpp>
+#include <openvino/op/parameter.hpp>
+#include <openvino/pass/constant_folding.hpp>
+
+#include "ngraph_functions/builders.hpp"
+
+namespace LayerTestsDefinitions {
+
+std::string EyeLayerTest::getTestCaseName(testing::TestParamInfo<EyeLayerTestParams> obj) {
+    EyeLayerTestParams params = obj.param;
+
+    std::string td;
+    std::vector<ov::Shape> input_shapes;
+    ElementType net_precision;
+    std::vector<LocalElementType> out_batch_shape;
+    std::vector<int> eye_par;
+    std::tie(input_shapes, out_batch_shape, eye_par, net_precision, td) = params;
+    std::ostringstream result;
+    result << "EyeTest_";
+    result << "IS=(";
+    for (const auto& shape : input_shapes) {
+        result << CommonTestUtils::partialShape2str({shape}) << "_";
+    }
+    result << ")";
+    result << "rowNum=" << eye_par[0] << "_";
+    result << "colNum=" << eye_par[1] << "_";
+    result << "diagShift=" << eye_par[2] << "_";
+    result << "batchShape=" << CommonTestUtils::vec2str(out_batch_shape) << "_";
+    result << net_precision << "_";
+    result << std::to_string(obj.index);
+    return result.str();
+}
+
+void EyeLayerTest::SetUp() {
+    std::vector<ov::Shape> input_shapes;
+    LocalElementType row_num, col_num;
+    LocalElementType shift;
+    std::vector<LocalElementType> out_batch_shape;
+    ElementType net_precision;
+    EyeLayerTestParams basicParamsSet = this->GetParam();
+
+    std::vector<int> eye_par;
+    std::tie(input_shapes, out_batch_shape, eye_par, net_precision, targetDevice) = basicParamsSet;
+    row_num = eye_par[0];
+    col_num = eye_par[1];
+    shift = eye_par[2];
+
+    std::shared_ptr<ngraph::op::v9::Eye> eye_operation;
+
+    auto rows_const = std::make_shared<ov::op::v0::Constant>(ngraph::element::i32, input_shapes[0], &row_num);
+    rows_const->set_friendly_name("rows");
+    auto cols_const = std::make_shared<ov::op::v0::Constant>(ngraph::element::i32, input_shapes[1], &col_num);
+    cols_const->set_friendly_name("cols");
+    auto diag_const = std::make_shared<ov::op::v0::Constant>(ngraph::element::i32, input_shapes[2], &shift);
+    diag_const->set_friendly_name("diagInd");
+
+    if (!out_batch_shape.empty() && out_batch_shape[0] != 0) {
+        auto batch_shape_par = std::make_shared<ov::op::v0::Constant>(ngraph::element::i32,
+                                                                      ov::Shape{out_batch_shape.size()},
+                                                                      out_batch_shape.data());
+        batch_shape_par->set_friendly_name("batchShape");
+        eye_operation =
+            std::make_shared<ngraph::op::v9::Eye>(rows_const, cols_const, diag_const, batch_shape_par, net_precision);
+    } else {
+        eye_operation = std::make_shared<ngraph::op::v9::Eye>(rows_const, cols_const, diag_const, net_precision);
+    }
+    // Without this call the eye operation will be calculated by CPU and substituted by Constant operator
+    ov::pass::disable_constant_folding(eye_operation);
+    ngraph::ResultVector results{std::make_shared<ngraph::opset1::Result>(eye_operation)};
+    function = std::make_shared<ngraph::Function>(results, ngraph::ParameterVector{}, "eye");
+}
+}  // namespace LayerTestsDefinitions