[GPU] Impelement ExperimentalDetectronPriorGridGenerator-6 gpu implementation (#11632)

* Implement experimental_detectron_prior_grid_generator kernel

* register experimental_detectron_prior_grid_generator operation

* implement single layer tests

* add unit tests to detectron

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

@@ -206,6 +206,7 @@ REGISTER_FACTORY(v5, Loop);
 REGISTER_FACTORY(v6, CTCGreedyDecoderSeqLen);
 REGISTER_FACTORY(v6, MVN);
 REGISTER_FACTORY(v6, GatherElements);
+REGISTER_FACTORY(v6, ExperimentalDetectronPriorGridGenerator);
 REGISTER_FACTORY(v6, ExperimentalDetectronROIFeatureExtractor);
 REGISTER_FACTORY(v6, ExperimentalDetectronTopKROIs)
 REGISTER_FACTORY(v6, ExperimentalDetectronGenerateProposalsSingleImage);

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

@@ -0,0 +1,57 @@
+// 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 Constructs experimental_detectron_prior_grid_generator primitive.
+struct experimental_detectron_prior_grid_generator
+    : public primitive_base<experimental_detectron_prior_grid_generator> {
+    CLDNN_DECLARE_PRIMITIVE(experimental_detectron_prior_grid_generator)
+
+    experimental_detectron_prior_grid_generator(const primitive_id& id,
+                                                const std::vector<primitive_id>& input,
+                                                bool flatten,
+                                                uint64_t h,
+                                                uint64_t w,
+                                                float stride_x,
+                                                float stride_y,
+                                                uint64_t featmap_height,
+                                                uint64_t featmap_width,
+                                                uint64_t image_height,
+                                                uint64_t image_width,
+                                                const primitive_id& ext_prim_id = {})
+        : primitive_base{id, input, ext_prim_id},
+          flatten{flatten},
+          h{h},
+          w{w},
+          stride_x{stride_x},
+          stride_y{stride_y},
+          featmap_height{featmap_height},
+          featmap_width{featmap_width},
+          image_height{image_height},
+          image_width{image_width} {}
+
+    bool flatten;
+    uint64_t h;
+    uint64_t w;
+    float stride_x;
+    float stride_y;
+    uint64_t featmap_height;
+    uint64_t featmap_width;
+    uint64_t image_height;
+    uint64_t image_width;
+};
+
+}  // namespace cldnn

src/plugins/intel_gpu/src/graph/experimental_detectron_prior_grid_generator.cpp

@@ -0,0 +1,51 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include <experimental_detectron_prior_grid_generator_inst.h>
+#include <primitive_type_base.h>
+
+#include "json_object.h"
+
+namespace cldnn {
+
+primitive_type_id experimental_detectron_prior_grid_generator::type_id() {
+    static primitive_type_base<experimental_detectron_prior_grid_generator> instance;
+    return &instance;
+}
+
+layout experimental_detectron_prior_grid_generator_inst::calc_output_layout(
+    const experimental_detectron_prior_grid_generator_node& node) {
+    const layout data_layout = node.input().get_output_layout();
+    auto desc = node.get_primitive();
+    if (desc->flatten) {
+        return layout(data_layout.data_type,
+                      format::bfyx,
+                      {static_cast<int>(desc->featmap_width * desc->featmap_height * data_layout.batch()), 4, 1, 1});
+    } else {
+        return layout(data_layout.data_type,
+                      format::bfyx,
+                      {static_cast<int>(desc->featmap_height),
+                       static_cast<int>(desc->featmap_width),
+                       4,
+                       static_cast<int>(data_layout.batch())});
+    }
+}
+
+std::string experimental_detectron_prior_grid_generator_inst::to_string(
+    experimental_detectron_prior_grid_generator_node const& node) {
+    auto node_info = node.desc_to_json();
+    json_composite experimental_detectron_prior_grid_generator_info;
+    experimental_detectron_prior_grid_generator_info.add("flatten", node.get_primitive()->flatten);
+    experimental_detectron_prior_grid_generator_info.add("h", node.get_primitive()->h);
+    experimental_detectron_prior_grid_generator_info.add("w", node.get_primitive()->w);
+    experimental_detectron_prior_grid_generator_info.add("stride_x", node.get_primitive()->stride_x);
+    experimental_detectron_prior_grid_generator_info.add("stride_y", node.get_primitive()->stride_y);
+    node_info->add("experimental_detectron_prior_grid_generator_info",
+                   experimental_detectron_prior_grid_generator_info);
+    std::stringstream primitive_description;
+    node_info->dump(primitive_description);
+    return primitive_description.str();
+}
+
+}  // namespace cldnn

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

@@ -0,0 +1,70 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include "intel_gpu/primitives/experimental_detectron_prior_grid_generator.hpp"
+
+#include <edpgg/experimental_detectron_prior_grid_generator_kernel_ref.h>
+#include <edpgg/experimental_detectron_prior_grid_generator_kernel_selector.h>
+#include <experimental_detectron_prior_grid_generator_inst.h>
+#include <kernel_selector_helper.h>
+
+#include <impls/implementation_map.hpp>
+#include <intel_gpu/runtime/error_handler.hpp>
+
+#include "primitive_base.hpp"
+
+namespace cldnn {
+namespace ocl {
+
+struct experimental_detectron_prior_grid_generator_impl
+    : typed_primitive_impl_ocl<experimental_detectron_prior_grid_generator> {
+    using parent = typed_primitive_impl_ocl<experimental_detectron_prior_grid_generator>;
+    using parent::parent;
+
+    std::unique_ptr<primitive_impl> clone() const override {
+        return make_unique<experimental_detectron_prior_grid_generator_impl>(*this);
+    }
+
+    static primitive_impl* create(const experimental_detectron_prior_grid_generator_node& arg) {
+        auto params = get_default_params<kernel_selector::experimental_detectron_prior_grid_generator_params>(arg);
+        auto primPtr = arg.get_primitive();
+        auto& prim = *primPtr;
+
+        params.flatten = prim.flatten;
+        params.layer_height = prim.h ? prim.h : prim.featmap_height;
+        params.layer_width = prim.w ? prim.w : prim.featmap_width;
+        params.step_x = prim.stride_x ? prim.stride_x : static_cast<float>(prim.image_width) / params.layer_width;
+        params.step_y = prim.stride_y ? prim.stride_y : static_cast<float>(prim.image_height) / params.layer_height;
+
+        auto optional_params =
+            get_default_optional_params<kernel_selector::experimental_detectron_prior_grid_generator_optional_params>(
+                arg.get_program());
+
+        auto& kernel_selector = kernel_selector::experimental_detectron_prior_grid_generator_instance();
+        auto best_kernels = kernel_selector.GetBestKernels(params, optional_params);
+
+        CLDNN_ERROR_BOOL(arg.id(),
+                         "Best_kernel.empty()",
+                         best_kernels.empty(),
+                         "Cannot find a proper kernel with this arguments");
+
+        return new experimental_detectron_prior_grid_generator_impl{arg, best_kernels.front()};
+    }
+};
+
+namespace detail {
+
+attach_experimental_detectron_prior_grid_generator_impl::attach_experimental_detectron_prior_grid_generator_impl() {
+    implementation_map<experimental_detectron_prior_grid_generator>::add(
+        impl_types::ocl,
+        experimental_detectron_prior_grid_generator_impl::create,
+        {
+            std::make_tuple(data_types::f16, format::bfyx),
+            std::make_tuple(data_types::f32, format::bfyx),
+        });
+}
+
+}  // namespace detail
+}  // namespace ocl
+}  // namespace cldnn

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

@@ -31,6 +31,7 @@ void register_implementations() {
     REGISTER_OCL(dft);
     REGISTER_OCL(batch_to_space);
     REGISTER_OCL(experimental_detectron_generate_proposals_single_image);
+    REGISTER_OCL(experimental_detectron_prior_grid_generator);
     REGISTER_OCL(experimental_detectron_roi_feature_extractor);
     REGISTER_OCL(experimental_detectron_topk_rois);
     REGISTER_OCL(eltwise);

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

@@ -20,6 +20,7 @@
 #include "intel_gpu/primitives/depth_to_space.hpp"
 #include "intel_gpu/primitives/detection_output.hpp"
 #include "intel_gpu/primitives/eltwise.hpp"
+#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/fully_connected.hpp"
@@ -103,6 +104,7 @@ REGISTER_OCL(depth_to_space);
 REGISTER_OCL(detection_output);
 REGISTER_OCL(dft);
 REGISTER_OCL(experimental_detectron_generate_proposals_single_image);
+REGISTER_OCL(experimental_detectron_prior_grid_generator);
 REGISTER_OCL(experimental_detectron_roi_feature_extractor);
 REGISTER_OCL(experimental_detectron_topk_rois);
 REGISTER_OCL(eltwise);

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

@@ -0,0 +1,38 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma once
+
+#include <intel_gpu/primitives/experimental_detectron_prior_grid_generator.hpp>
+
+#include "primitive_inst.h"
+
+namespace cldnn {
+template <>
+struct typed_program_node<experimental_detectron_prior_grid_generator>
+    : public typed_program_node_base<experimental_detectron_prior_grid_generator> {
+    using typed_program_node_base::typed_program_node_base;
+
+    program_node& input(std::size_t i = 0) const {
+        return get_dependency(i);
+    }
+};
+using experimental_detectron_prior_grid_generator_node =
+    typed_program_node<experimental_detectron_prior_grid_generator>;
+
+template <>
+class typed_primitive_inst<experimental_detectron_prior_grid_generator>
+    : public typed_primitive_inst_base<experimental_detectron_prior_grid_generator> {
+public:
+    using typed_primitive_inst_base::typed_primitive_inst_base;
+
+    static layout calc_output_layout(experimental_detectron_prior_grid_generator_node const& node);
+
+    static std::string to_string(experimental_detectron_prior_grid_generator_node const& node);
+};
+
+using experimental_detectron_prior_grid_generator_inst =
+    typed_primitive_inst<experimental_detectron_prior_grid_generator>;
+
+}  // namespace cldnn

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

@@ -78,6 +78,7 @@ enum class KernelType {
     NON_MAX_SUPPRESSION,
     DETECTION_OUTPUT,
     EXPERIMENTAL_DETECTRON_GENERATE_PROPOSALS_SINGLE_IMAGE,
+    EXPERIMENTAL_DETECTRON_PRIOR_GRID_GENERATOR,
     EXPERIMENTAL_DETECTRON_ROI_FEATURE_EXTRACTOR,
     EXPERIMENTAL_DETECTRON_TOPK_ROIS,
     CONVERT_COLOR,

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/edpgg/experimental_detectron_prior_grid_generator_kernel_ref.cpp

@@ -0,0 +1,77 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include "experimental_detectron_prior_grid_generator_kernel_ref.h"
+
+#include <kernel_selector_utils.h>
+
+namespace kernel_selector {
+
+namespace {
+
+CommonDispatchData SetDefault(const experimental_detectron_prior_grid_generator_params &params) {
+    CommonDispatchData dispatchData;
+
+    dispatchData.gws = { params.layer_height, params.layer_width, params.inputs[0].Batch().v };
+    dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, params.engineInfo);
+
+    return dispatchData;
+}
+
+}  // namespace
+
+KernelsData ExperimentalDetectronPriorGridGeneratorKernelRef::GetKernelsData(const Params &params, const optional_params &options) const {
+    KernelsData kernels_data;
+    if (!Validate(params, options))
+        return kernels_data;
+    kernels_data.push_back(KernelData::Default<experimental_detectron_prior_grid_generator_params>(params));
+    KernelData &kernel_data = kernels_data.front();
+    auto &derived_params = dynamic_cast<experimental_detectron_prior_grid_generator_params&>(*kernel_data.params.get());
+    auto dispatch_data = SetDefault(derived_params);
+    auto entry_point = GetEntryPoint(kernelName, derived_params.layerID, params, options);
+    auto jit_constants = GetJitConstants(derived_params);
+    auto jit = CreateJit(kernelName, jit_constants, entry_point);
+    auto &clKernelData = kernel_data.kernels[0];
+    FillCLKernelData(clKernelData, dispatch_data, params.engineInfo, kernelName, jit, entry_point);
+    return kernels_data;
+}
+
+KernelsPriority ExperimentalDetectronPriorGridGeneratorKernelRef::GetKernelsPriority(const Params& /*params*/, const optional_params& /*options*/) const {
+    return DONT_USE_IF_HAVE_SOMETHING_ELSE;
+}
+
+ParamsKey ExperimentalDetectronPriorGridGeneratorKernelRef::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.EnableBatching();
+    return k;
+}
+
+bool ExperimentalDetectronPriorGridGeneratorKernelRef::Validate(const Params &p, const optional_params &o) const {
+    if (p.GetType() != KernelType::EXPERIMENTAL_DETECTRON_PRIOR_GRID_GENERATOR || o.GetType() != KernelType::EXPERIMENTAL_DETECTRON_PRIOR_GRID_GENERATOR)
+        return false;
+
+    auto &params = dynamic_cast<const experimental_detectron_prior_grid_generator_params&>(p);
+    if (params.inputs.size() != 1)
+        return false;
+
+    return true;
+}
+
+JitConstants ExperimentalDetectronPriorGridGeneratorKernelRef::GetJitConstants(const experimental_detectron_prior_grid_generator_params& params) const {
+    auto jit_constants = MakeBaseParamsJitConstants(params);
+    jit_constants.AddConstant(MakeJitConstant("LAYER_WIDTH", params.layer_width));
+    jit_constants.AddConstant(MakeJitConstant("STEP_X", params.step_x));
+    jit_constants.AddConstant(MakeJitConstant("STEP_Y", params.step_y));
+    if (params.flatten) {
+        jit_constants.AddConstant(MakeJitConstant("FLATTEN", 1));
+    }
+    return jit_constants;
+}
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/edpgg/experimental_detectron_prior_grid_generator_kernel_ref.h

@@ -0,0 +1,53 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma once
+
+#include "kernel_base_opencl.h"
+
+namespace kernel_selector {
+
+/**
+ * ExperimentalDetectronPriorGridGenerator reference kernel parameters
+ */
+struct experimental_detectron_prior_grid_generator_params : public base_params {
+    experimental_detectron_prior_grid_generator_params()
+        : base_params{KernelType::EXPERIMENTAL_DETECTRON_PRIOR_GRID_GENERATOR} {}
+
+    bool flatten = false;
+    size_t layer_height = 0;
+    size_t layer_width = 0;
+    float step_x = 0;
+    float step_y = 0;
+};
+
+/**
+ * ExperimentalDetectronPriorGridGenerator reference kernel optional parameters
+ */
+struct experimental_detectron_prior_grid_generator_optional_params : optional_params {
+    experimental_detectron_prior_grid_generator_optional_params()
+        : optional_params{KernelType::EXPERIMENTAL_DETECTRON_PRIOR_GRID_GENERATOR} {}
+};
+
+/**
+ * Reference kernel for ExperimentalDetectronPriorGridGenerator
+ */
+class ExperimentalDetectronPriorGridGeneratorKernelRef : public KernelBaseOpenCL {
+public:
+    ExperimentalDetectronPriorGridGeneratorKernelRef()
+        : KernelBaseOpenCL{"experimental_detectron_prior_grid_generator_ref"} {}
+
+private:
+    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;
+
+    JitConstants GetJitConstants(const experimental_detectron_prior_grid_generator_params &params) const;
+};
+
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/edpgg/experimental_detectron_prior_grid_generator_kernel_selector.cpp

@@ -0,0 +1,28 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include "experimental_detectron_prior_grid_generator_kernel_selector.h"
+#include "experimental_detectron_prior_grid_generator_kernel_ref.h"
+
+namespace kernel_selector {
+namespace {
+
+class experimental_detectron_prior_grid_generator_kernel_selector : public kernel_selector_base {
+    KernelsData GetBestKernels(const Params &params, const optional_params &options) const override {
+        return GetNaiveBestKernel(params, options, KernelType::EXPERIMENTAL_DETECTRON_PRIOR_GRID_GENERATOR);
+    }
+public:
+    experimental_detectron_prior_grid_generator_kernel_selector() {
+        Attach<ExperimentalDetectronPriorGridGeneratorKernelRef>();
+    }
+};
+
+}  // namespace
+
+kernel_selector_base& experimental_detectron_prior_grid_generator_instance() {
+    static experimental_detectron_prior_grid_generator_kernel_selector instance;
+    return instance;
+}
+
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/edpgg/experimental_detectron_prior_grid_generator_kernel_selector.h

@@ -0,0 +1,16 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma once
+
+#include <kernel_selector.h>
+
+namespace kernel_selector {
+
+/*
+ * ExperimentalDetectronPriorGridGenerator kernel selector
+ */
+kernel_selector_base& experimental_detectron_prior_grid_generator_instance();
+
+}  // namespace kernel_selector

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

@@ -0,0 +1,30 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#if OUTPUT_TYPE_SIZE == 2 //f16
+#define HALF_ONE 0.5h
+#else
+#define HALF_ONE 0.5f
+#endif
+
+KERNEL (experimental_detectron_prior_grid_generator_ref)(const __global INPUT0_TYPE *input, __global OUTPUT_TYPE *output)
+{
+    const size_t y = get_global_id(0);
+    const size_t x = get_global_id(1);
+    const size_t prior = get_global_id(2);
+    const ptrdiff_t inputOffset = INPUT0_GET_INDEX(prior, 0, 0, 0);
+    const OUTPUT_TYPE stepX = STEP_X * ( x + HALF_ONE );
+    const OUTPUT_TYPE stepY = STEP_Y * ( y + HALF_ONE );
+#ifdef FLATTEN
+    const ptrdiff_t outputOffset = OUTPUT_GET_INDEX(LAYER_WIDTH * INPUT0_BATCH_NUM * y + INPUT0_BATCH_NUM * x + prior, 0, 0, 0);
+#else
+    const ptrdiff_t outputOffset = OUTPUT_OFFSET + OUTPUT_PITCHES[1] * ( LAYER_WIDTH * INPUT0_BATCH_NUM * y + INPUT0_BATCH_NUM * x + prior);
+#endif
+    output[outputOffset] = input[inputOffset] + stepX;
+    output[outputOffset + 1] =  input[inputOffset + 1] + stepY;
+    output[outputOffset + 2] = input[inputOffset + 2] + stepX;
+    output[outputOffset + 3] = input[inputOffset + 3] + stepY;
+}
+
+#undef HALF_ONE

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

@@ -0,0 +1,54 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include <intel_gpu/plugin/common_utils.hpp>
+#include <intel_gpu/plugin/program.hpp>
+#include <intel_gpu/primitives/experimental_detectron_prior_grid_generator.hpp>
+#include <ngraph/op/experimental_detectron_prior_grid_generator.hpp>
+
+namespace ov {
+namespace runtime {
+namespace intel_gpu {
+namespace {
+cldnn::tensor mkTensor(const ov::Shape& shape) {
+    if (shape.size() == 4)
+        return cldnn::tensor{cldnn::batch(shape[0]), cldnn::feature(shape[1]), cldnn::spatial(4, shape[2])};
+    else
+        return cldnn::tensor{cldnn::spatial(4, shape[0])};
+}
+
+static void CreateExperimentalDetectronPriorGridGeneratorOp(
+    Program& p,
+    const std::shared_ptr<ngraph::op::v6::ExperimentalDetectronPriorGridGenerator>& op) {
+    p.ValidateInputs(op, {3});
+    cldnn::tensor outTensor = mkTensor(op->get_output_shape(0));
+    auto outDataType = DataTypeFromPrecision(op->get_output_element_type(0));
+    cldnn::layout outLayout{outDataType, cldnn::format::bfyx, outTensor};
+    auto& attrs = op->get_attrs();
+    auto& featmap_shape = op->get_input_shape(1);
+    auto& image_shape = op->get_input_shape(2);
+    auto inputPrimitives = p.GetInputPrimitiveIDs(op);
+    inputPrimitives.resize(1);  // only priors is read
+    cldnn::experimental_detectron_prior_grid_generator prim{layer_type_name_ID(op),
+                                                            inputPrimitives,
+                                                            attrs.flatten,
+                                                            static_cast<uint64_t>(attrs.h),
+                                                            static_cast<uint64_t>(attrs.w),
+                                                            attrs.stride_x,
+                                                            attrs.stride_y,
+                                                            featmap_shape[2],
+                                                            featmap_shape[3],
+                                                            image_shape[2],
+                                                            image_shape[3],
+                                                            op->get_friendly_name()};
+    p.AddPrimitive(prim);
+    p.AddPrimitiveToProfiler(op);
+}
+}  // namespace
+
+REGISTER_FACTORY_IMPL(v6, ExperimentalDetectronPriorGridGenerator);
+
+}  // namespace intel_gpu
+}  // namespace runtime
+}  // namespace ov

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

@@ -0,0 +1,230 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include <cstddef>
+#include <intel_gpu/primitives/activation.hpp>
+#include <intel_gpu/primitives/experimental_detectron_prior_grid_generator.hpp>
+#include <intel_gpu/primitives/input_layout.hpp>
+#include <string>
+
+#include "test_utils.h"
+
+using namespace cldnn;
+using namespace ::tests;
+
+template <typename T>
+struct ExperimentalDetectronPriorGridGeneratorParams {
+    std::vector<T> priors;
+    tensor priorsTensor;
+    int h;
+    int w;
+    float strideX;
+    float strideY;
+    bool flatten;
+    std::pair<int, int> featureShape;
+    std::pair<int, int> imageShape;
+    tensor outputTensor;
+    std::vector<T> expectedOutput;
+};
+
+template <typename T>
+struct experimental_detectron_prior_grid_generator_test
+    : public ::testing::TestWithParam<ExperimentalDetectronPriorGridGeneratorParams<T>> {
+public:
+    void test() {
+        auto data_type = type_to_data_type<T>::value;
+        ExperimentalDetectronPriorGridGeneratorParams<T> params =
+            testing::TestWithParam<ExperimentalDetectronPriorGridGeneratorParams<T>>::GetParam();
+        auto& engine = get_test_engine();
+
+        auto prior_input = engine.allocate_memory({data_type, format::bfyx, params.priorsTensor});
+
+        set_values(prior_input, params.priors);
+
+        const std::string priors_id = "priors";
+        const std::string experimental_detectron_prior_grid_generator_id =
+            "experimental_detectron_prior_grid_generator";
+        topology topology;
+        topology.add(input_layout(priors_id, prior_input->get_layout()));
+
+        cldnn::layout outLayout{data_type, cldnn::format::bfyx, params.outputTensor};
+        topology.add(experimental_detectron_prior_grid_generator(experimental_detectron_prior_grid_generator_id,
+                                                                 {priors_id},
+                                                                 params.flatten,
+                                                                 params.h,
+                                                                 params.w,
+                                                                 params.strideX,
+                                                                 params.strideY,
+                                                                 params.featureShape.first,
+                                                                 params.featureShape.second,
+                                                                 params.imageShape.first,
+                                                                 params.imageShape.second));
+
+        network network(engine, topology);
+
+        network.set_input_data(priors_id, prior_input);
+
+        auto result = network.execute();
+
+        auto out_mem = result.at(experimental_detectron_prior_grid_generator_id).get_memory();
+        cldnn::mem_lock<T> out_ptr(out_mem, get_test_stream());
+
+        ASSERT_EQ(params.outputTensor.count(), out_ptr.size());
+        for (size_t i = 0; i < params.expectedOutput.size(); ++i) {
+            EXPECT_NEAR(params.expectedOutput[i], out_ptr[i], 0.0001) << "at i = " << i;
+        }
+    }
+};
+
+template <typename T>
+std::vector<T> getValues(const std::vector<float>& values) {
+    std::vector<T> result(values.begin(), values.end());
+    return result;
+}
+
+template <typename T>
+std::vector<ExperimentalDetectronPriorGridGeneratorParams<T>> generateExperimentalPGGParams() {
+    std::vector<ExperimentalDetectronPriorGridGeneratorParams<T>> experimentalPGGParams{
+        {getValues<T>({-24.5, -12.5, 24.5, 12.5, -16.5, -16.5, 16.5, 16.5, -12.5, -24.5, 12.5, 24.5}),
+         tensor(batch(3), feature(4)),
+         0,
+         0,
+         4.0,
+         4.0,
+         true,
+         {4, 5},
+         {100, 200},
+         tensor(1, 1, 4, 60),
+         getValues<T>(
+             {-22.5, -10.5, 26.5, 14.5, -14.5, -14.5, 18.5, 18.5, -10.5, -22.5, 14.5, 26.5, -18.5, -10.5, 30.5, 14.5,
+              -10.5, -14.5, 22.5, 18.5, -6.5,  -22.5, 18.5, 26.5, -14.5, -10.5, 34.5, 14.5, -6.5,  -14.5, 26.5, 18.5,
+              -2.5,  -22.5, 22.5, 26.5, -10.5, -10.5, 38.5, 14.5, -2.5,  -14.5, 30.5, 18.5, 1.5,   -22.5, 26.5, 26.5,
+              -6.5,  -10.5, 42.5, 14.5, 1.5,   -14.5, 34.5, 18.5, 5.5,   -22.5, 30.5, 26.5, -22.5, -6.5,  26.5, 18.5,
+              -14.5, -10.5, 18.5, 22.5, -10.5, -18.5, 14.5, 30.5, -18.5, -6.5,  30.5, 18.5, -10.5, -10.5, 22.5, 22.5,
+              -6.5,  -18.5, 18.5, 30.5, -14.5, -6.5,  34.5, 18.5, -6.5,  -10.5, 26.5, 22.5, -2.5,  -18.5, 22.5, 30.5,
+              -10.5, -6.5,  38.5, 18.5, -2.5,  -10.5, 30.5, 22.5, 1.5,   -18.5, 26.5, 30.5, -6.5,  -6.5,  42.5, 18.5,
+              1.5,   -10.5, 34.5, 22.5, 5.5,   -18.5, 30.5, 30.5, -22.5, -2.5,  26.5, 22.5, -14.5, -6.5,  18.5, 26.5,
+              -10.5, -14.5, 14.5, 34.5, -18.5, -2.5,  30.5, 22.5, -10.5, -6.5,  22.5, 26.5, -6.5,  -14.5, 18.5, 34.5,
+              -14.5, -2.5,  34.5, 22.5, -6.5,  -6.5,  26.5, 26.5, -2.5,  -14.5, 22.5, 34.5, -10.5, -2.5,  38.5, 22.5,
+              -2.5,  -6.5,  30.5, 26.5, 1.5,   -14.5, 26.5, 34.5, -6.5,  -2.5,  42.5, 22.5, 1.5,   -6.5,  34.5, 26.5,
+              5.5,   -14.5, 30.5, 34.5, -22.5, 1.5,   26.5, 26.5, -14.5, -2.5,  18.5, 30.5, -10.5, -10.5, 14.5, 38.5,
+              -18.5, 1.5,   30.5, 26.5, -10.5, -2.5,  22.5, 30.5, -6.5,  -10.5, 18.5, 38.5, -14.5, 1.5,   34.5, 26.5,
+              -6.5,  -2.5,  26.5, 30.5, -2.5,  -10.5, 22.5, 38.5, -10.5, 1.5,   38.5, 26.5, -2.5,  -2.5,  30.5, 30.5,
+              1.5,   -10.5, 26.5, 38.5, -6.5,  1.5,   42.5, 26.5, 1.5,   -2.5,  34.5, 30.5, 5.5,   -10.5, 30.5, 38.5})},
+        {getValues<T>({-44.5, -24.5, 44.5, 24.5, -32.5, -32.5, 32.5, 32.5, -24.5, -44.5, 24.5, 44.5}),
+         tensor(batch(3), feature(4)),
+         0,
+         0,
+         8.0,
+         8.0,
+         false,
+         {3, 7},
+         {100, 200},
+         tensor(3, 7, 4, 3),
+         getValues<T>(
+             {-40.5, -20.5, 48.5, 28.5, -28.5, -28.5, 36.5, 36.5, -20.5, -40.5, 28.5, 48.5, -32.5, -20.5, 56.5, 28.5,
+              -20.5, -28.5, 44.5, 36.5, -12.5, -40.5, 36.5, 48.5, -24.5, -20.5, 64.5, 28.5, -12.5, -28.5, 52.5, 36.5,
+              -4.5,  -40.5, 44.5, 48.5, -16.5, -20.5, 72.5, 28.5, -4.5,  -28.5, 60.5, 36.5, 3.5,   -40.5, 52.5, 48.5,
+              -8.5,  -20.5, 80.5, 28.5, 3.5,   -28.5, 68.5, 36.5, 11.5,  -40.5, 60.5, 48.5, -0.5,  -20.5, 88.5, 28.5,
+              11.5,  -28.5, 76.5, 36.5, 19.5,  -40.5, 68.5, 48.5, 7.5,   -20.5, 96.5, 28.5, 19.5,  -28.5, 84.5, 36.5,
+              27.5,  -40.5, 76.5, 48.5, -40.5, -12.5, 48.5, 36.5, -28.5, -20.5, 36.5, 44.5, -20.5, -32.5, 28.5, 56.5,
+              -32.5, -12.5, 56.5, 36.5, -20.5, -20.5, 44.5, 44.5, -12.5, -32.5, 36.5, 56.5, -24.5, -12.5, 64.5, 36.5,
+              -12.5, -20.5, 52.5, 44.5, -4.5,  -32.5, 44.5, 56.5, -16.5, -12.5, 72.5, 36.5, -4.5,  -20.5, 60.5, 44.5,
+              3.5,   -32.5, 52.5, 56.5, -8.5,  -12.5, 80.5, 36.5, 3.5,   -20.5, 68.5, 44.5, 11.5,  -32.5, 60.5, 56.5,
+              -0.5,  -12.5, 88.5, 36.5, 11.5,  -20.5, 76.5, 44.5, 19.5,  -32.5, 68.5, 56.5, 7.5,   -12.5, 96.5, 36.5,
+              19.5,  -20.5, 84.5, 44.5, 27.5,  -32.5, 76.5, 56.5, -40.5, -4.5,  48.5, 44.5, -28.5, -12.5, 36.5, 52.5,
+              -20.5, -24.5, 28.5, 64.5, -32.5, -4.5,  56.5, 44.5, -20.5, -12.5, 44.5, 52.5, -12.5, -24.5, 36.5, 64.5,
+              -24.5, -4.5,  64.5, 44.5, -12.5, -12.5, 52.5, 52.5, -4.5,  -24.5, 44.5, 64.5, -16.5, -4.5,  72.5, 44.5,
+              -4.5,  -12.5, 60.5, 52.5, 3.5,   -24.5, 52.5, 64.5, -8.5,  -4.5,  80.5, 44.5, 3.5,   -12.5, 68.5, 52.5,
+              11.5,  -24.5, 60.5, 64.5, -0.5,  -4.5,  88.5, 44.5, 11.5,  -12.5, 76.5, 52.5, 19.5,  -24.5, 68.5, 64.5,
+              7.5,   -4.5,  96.5, 44.5, 19.5,  -12.5, 84.5, 52.5, 27.5,  -24.5, 76.5, 64.5})
+
+        },
+        {getValues<T>({-364.5, -184.5, 364.5, 184.5, -256.5, -256.5, 256.5, 256.5, -180.5, -360.5, 180.5, 360.5}),
+         tensor(batch(3), feature(4)),
+         3,
+         6,
+         64.0,
+         64.0,
+         true,
+         {100, 100},
+         {100, 200},
+         tensor(1, 1, 4, 30000),
+         getValues<T>({-332.5, -152.5, 396.5, 216.5, -224.5, -224.5, 288.5, 288.5, -148.5, -328.5, 212.5, 392.5,
+                       -268.5, -152.5, 460.5, 216.5, -160.5, -224.5, 352.5, 288.5, -84.5,  -328.5, 276.5, 392.5,
+                       -204.5, -152.5, 524.5, 216.5, -96.5,  -224.5, 416.5, 288.5, -20.5,  -328.5, 340.5, 392.5,
+                       -140.5, -152.5, 588.5, 216.5, -32.5,  -224.5, 480.5, 288.5, 43.5,   -328.5, 404.5, 392.5,
+                       -76.5,  -152.5, 652.5, 216.5, 31.5,   -224.5, 544.5, 288.5, 107.5,  -328.5, 468.5, 392.5,
+                       -12.5,  -152.5, 716.5, 216.5, 95.5,   -224.5, 608.5, 288.5, 171.5,  -328.5, 532.5, 392.5,
+                       -332.5, -88.5,  396.5, 280.5, -224.5, -160.5, 288.5, 352.5, -148.5, -264.5, 212.5, 456.5,
+                       -268.5, -88.5,  460.5, 280.5, -160.5, -160.5, 352.5, 352.5, -84.5,  -264.5, 276.5, 456.5,
+                       -204.5, -88.5,  524.5, 280.5, -96.5,  -160.5, 416.5, 352.5, -20.5,  -264.5, 340.5, 456.5,
+                       -140.5, -88.5,  588.5, 280.5, -32.5,  -160.5, 480.5, 352.5, 43.5,   -264.5, 404.5, 456.5,
+                       -76.5,  -88.5,  652.5, 280.5, 31.5,   -160.5, 544.5, 352.5, 107.5,  -264.5, 468.5, 456.5,
+                       -12.5,  -88.5,  716.5, 280.5, 95.5,   -160.5, 608.5, 352.5, 171.5,  -264.5, 532.5, 456.5,
+                       -332.5, -24.5,  396.5, 344.5, -224.5, -96.5,  288.5, 416.5, -148.5, -200.5, 212.5, 520.5,
+                       -268.5, -24.5,  460.5, 344.5, -160.5, -96.5,  352.5, 416.5, -84.5,  -200.5, 276.5, 520.5,
+                       -204.5, -24.5,  524.5, 344.5, -96.5,  -96.5,  416.5, 416.5, -20.5,  -200.5, 340.5, 520.5,
+                       -140.5, -24.5,  588.5, 344.5, -32.5,  -96.5,  480.5, 416.5, 43.5,   -200.5, 404.5, 520.5,
+                       -76.5,  -24.5,  652.5, 344.5, 31.5,   -96.5,  544.5, 416.5, 107.5,  -200.5, 468.5, 520.5,
+                       -12.5,  -24.5,  716.5, 344.5, 95.5,   -96.5,  608.5, 416.5, 171.5,  -200.5, 532.5, 520.5})},
+        {getValues<T>({-180.5, -88.5, 180.5, 88.5, -128.5, -128.5, 128.5, 128.5, -92.5, -184.5, 92.5, 184.5}),
+         tensor(batch(3), feature(4)),
+         5,
+         3,
+         32.0,
+         32.0,
+         false,
+         {100, 100},
+         {100, 200},
+         tensor(100, 100, 4, 3),
+         getValues<T>({-164.5, -72.5, 196.5, 104.5, -112.5, -112.5, 144.5, 144.5, -76.5, -168.5, 108.5, 200.5,
+                       -132.5, -72.5, 228.5, 104.5, -80.5,  -112.5, 176.5, 144.5, -44.5, -168.5, 140.5, 200.5,
+                       -100.5, -72.5, 260.5, 104.5, -48.5,  -112.5, 208.5, 144.5, -12.5, -168.5, 172.5, 200.5,
+                       -164.5, -40.5, 196.5, 136.5, -112.5, -80.5,  144.5, 176.5, -76.5, -136.5, 108.5, 232.5,
+                       -132.5, -40.5, 228.5, 136.5, -80.5,  -80.5,  176.5, 176.5, -44.5, -136.5, 140.5, 232.5,
+                       -100.5, -40.5, 260.5, 136.5, -48.5,  -80.5,  208.5, 176.5, -12.5, -136.5, 172.5, 232.5,
+                       -164.5, -8.5,  196.5, 168.5, -112.5, -48.5,  144.5, 208.5, -76.5, -104.5, 108.5, 264.5,
+                       -132.5, -8.5,  228.5, 168.5, -80.5,  -48.5,  176.5, 208.5, -44.5, -104.5, 140.5, 264.5,
+                       -100.5, -8.5,  260.5, 168.5, -48.5,  -48.5,  208.5, 208.5, -12.5, -104.5, 172.5, 264.5,
+                       -164.5, 23.5,  196.5, 200.5, -112.5, -16.5,  144.5, 240.5, -76.5, -72.5,  108.5, 296.5,
+                       -132.5, 23.5,  228.5, 200.5, -80.5,  -16.5,  176.5, 240.5, -44.5, -72.5,  140.5, 296.5,
+                       -100.5, 23.5,  260.5, 200.5, -48.5,  -16.5,  208.5, 240.5, -12.5, -72.5,  172.5, 296.5,
+                       -164.5, 55.5,  196.5, 232.5, -112.5, 15.5,   144.5, 272.5, -76.5, -40.5,  108.5, 328.5,
+                       -132.5, 55.5,  228.5, 232.5, -80.5,  15.5,   176.5, 272.5, -44.5, -40.5,  140.5, 328.5,
+                       -100.5, 55.5,  260.5, 232.5, -48.5,  15.5,   208.5, 272.5, -12.5, -40.5,  172.5, 328.5})}};
+    return experimentalPGGParams;
+}
+
+struct PrintToStringParamName {
+    template <class T>
+    std::string operator()(const testing::TestParamInfo<ExperimentalDetectronPriorGridGeneratorParams<T>>& param) {
+        std::stringstream buf;
+        buf << " priors tensor " << param.param.priorsTensor.to_string() << " h " << param.param.h << " w "
+            << param.param.w << " strideX " << param.param.strideX << " strideY " << param.param.strideY << " flatten "
+            << param.param.flatten;
+        return buf.str();
+    }
+};
+
+using experimental_detectron_prior_grid_generator_test_f32 = experimental_detectron_prior_grid_generator_test<float>;
+using experimental_detectron_prior_grid_generator_test_f16 = experimental_detectron_prior_grid_generator_test<half_t>;
+
+TEST_P(experimental_detectron_prior_grid_generator_test_f32, experimental_detectron_prior_grid_generator_test_f32) {
+    ASSERT_NO_FATAL_FAILURE(test());
+}
+
+TEST_P(experimental_detectron_prior_grid_generator_test_f16, experimental_detectron_prior_grid_generator_test_f16) {
+    ASSERT_NO_FATAL_FAILURE(test());
+}
+
+INSTANTIATE_TEST_SUITE_P(smoke_experimental_detectron_prior_grid_generator_test_f32,
+                         experimental_detectron_prior_grid_generator_test_f32,
+                         ::testing::ValuesIn(generateExperimentalPGGParams<float>()),
+                         PrintToStringParamName());
+
+INSTANTIATE_TEST_SUITE_P(smoke_experimental_detectron_prior_grid_generator_test_f16,
+                         experimental_detectron_prior_grid_generator_test_f16,
+                         ::testing::ValuesIn(generateExperimentalPGGParams<half_t>()),
+                         PrintToStringParamName());

src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/experimental_detectron_prior_grid_generator.cpp

@@ -0,0 +1,74 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include <common_test_utils/data_utils.hpp>
+#include <common_test_utils/ov_tensor_utils.hpp>
+#include <single_layer_tests/experimental_detectron_prior_grid_generator.hpp>
+
+namespace {
+
+const std::initializer_list<ov::test::subgraph::ExperimentalDetectronPriorGridGeneratorTestParam> params{
+    // flatten = true (output tensor is 2D)
+    {{true, 0, 0, 4.0f, 4.0f},
+     ov::test::static_shapes_to_test_representation({{3, 4}, {1, 16, 4, 5}, {1, 3, 100, 200}})},
+    // task #72587
+    //{
+    //        {true, 3, 6, 64.0f, 64.0f},
+    //        ov::test::static_shapes_to_test_representation({{3, 4}, {1, 16, 100, 100}, {1, 3, 100, 200}})
+    //},
+    //// flatten = false (output tensor is 4D)
+    {{false, 0, 0, 8.0f, 8.0f},
+     ov::test::static_shapes_to_test_representation({{3, 4}, {1, 16, 3, 7}, {1, 3, 100, 200}})},
+    // task #72587
+    //{
+    //        {false, 5, 3, 32.0f, 32.0f},
+    //        ov::test::static_shapes_to_test_representation({{3, 4}, {1, 16, 100, 100}, {1, 3, 100, 200}})
+    //},
+};
+
+template <typename T>
+std::vector<std::pair<std::string, std::vector<ov::Tensor>>> getInputTensors() {
+    std::vector<std::pair<std::string, std::vector<ov::Tensor>>> tensors{
+        {"test#1",
+         {ov::test::utils::create_tensor<T>(
+             ov::element::from<T>(),
+             {3, 4},
+             std::vector<T>{-24.5, -12.5, 24.5, 12.5, -16.5, -16.5, 16.5, 16.5, -12.5, -24.5, 12.5, 24.5})}},
+        {"test#2",
+         {ov::test::utils::create_tensor<T>(
+             ov::element::from<T>(),
+             {3, 4},
+             std::vector<T>{-44.5, -24.5, 44.5, 24.5, -32.5, -32.5, 32.5, 32.5, -24.5, -44.5, 24.5, 44.5})}},
+        {"test#3",
+         {ov::test::utils::create_tensor<T>(
+             ov::element::from<T>(),
+             {3, 4},
+             std::
+                 vector<T>{-364.5, -184.5, 364.5, 184.5, -256.5, -256.5, 256.5, 256.5, -180.5, -360.5, 180.5, 360.5})}},
+        {"test#4",
+         {ov::test::utils::create_tensor<T>(
+             ov::element::from<T>(),
+             {3, 4},
+             std::vector<T>{-180.5, -88.5, 180.5, 88.5, -128.5, -128.5, 128.5, 128.5, -92.5, -184.5, 92.5, 184.5})}}};
+    return tensors;
+}
+
+using ov::test::subgraph::ExperimentalDetectronPriorGridGeneratorLayerTest;
+
+INSTANTIATE_TEST_SUITE_P(smoke_ExperimentalDetectronPriorGridGenerator_f32,
+                         ExperimentalDetectronPriorGridGeneratorLayerTest,
+                         testing::Combine(testing::ValuesIn(params),
+                                          testing::ValuesIn(getInputTensors<float>()),
+                                          testing::ValuesIn({ov::element::Type_t::f32}),
+                                          testing::Values(CommonTestUtils::DEVICE_GPU)),
+                         ExperimentalDetectronPriorGridGeneratorLayerTest::getTestCaseName);
+
+INSTANTIATE_TEST_SUITE_P(smoke_ExperimentalDetectronPriorGridGenerator_f16,
+                         ExperimentalDetectronPriorGridGeneratorLayerTest,
+                         testing::Combine(testing::ValuesIn(params),
+                                          testing::ValuesIn(getInputTensors<ov::float16>()),
+                                          testing::ValuesIn({ov::element::Type_t::f16}),
+                                          testing::Values(CommonTestUtils::DEVICE_GPU)),
+                         ExperimentalDetectronPriorGridGeneratorLayerTest::getTestCaseName);
+}  // namespace

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

@@ -70,8 +70,19 @@ void ExperimentalDetectronPriorGridGeneratorLayerTest::SetUp() {
             "ExperimentalDetectronPriorGridGenerator");
 }
 
+namespace {
+template<typename T>
+ov::runtime::Tensor generateTensorByShape(const Shape &shape) {
+    return ov::test::utils::create_tensor<T>(
+            ov::element::from<T>(),
+            shape,
+            std::vector<T>(0., shape_size(shape)));
+}
+}
+
 void ExperimentalDetectronPriorGridGeneratorLayerTest::generate_inputs(const std::vector<ngraph::Shape>& targetInputStaticShapes) {
     auto inputTensors = std::get<1>(GetParam());
+    auto netPrecision = std::get<2>(GetParam());
 
     inputs.clear();
     const auto& funcInputs = function->inputs();
@@ -84,12 +95,11 @@ void ExperimentalDetectronPriorGridGeneratorLayerTest::generate_inputs(const std
 
         inputs.insert({funcInputs[i].get_node_shared_ptr(), inputTensors.second[i]});
     }
-
     for (auto j = i; j < funcInputs.size(); ++j) {
-        ov::runtime::Tensor inputTensor = ov::test::utils::create_tensor<float>(
-            ov::element::f32,
-            targetInputStaticShapes[j],
-            std::vector<float>(0.f, shape_size(targetInputStaticShapes[j])));
+        ov::runtime::Tensor inputTensor = (netPrecision == element::f16)
+                                          ? generateTensorByShape<ov::float16>(targetInputStaticShapes[j])
+                                          : generateTensorByShape<float>(
+                        targetInputStaticShapes[j]);
 
         inputs.insert({funcInputs[j].get_node_shared_ptr(), inputTensor});
     }