diff --git a/src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp b/src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp index a6f6f9463cb..850c23644b5 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp @@ -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); diff --git a/src/plugins/intel_gpu/include/intel_gpu/primitives/experimental_detectron_prior_grid_generator.hpp b/src/plugins/intel_gpu/include/intel_gpu/primitives/experimental_detectron_prior_grid_generator.hpp new file mode 100644 index 00000000000..2269b2489c9 --- /dev/null +++ b/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 { + CLDNN_DECLARE_PRIMITIVE(experimental_detectron_prior_grid_generator) + + experimental_detectron_prior_grid_generator(const primitive_id& id, + const std::vector& 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 diff --git a/src/plugins/intel_gpu/src/graph/experimental_detectron_prior_grid_generator.cpp b/src/plugins/intel_gpu/src/graph/experimental_detectron_prior_grid_generator.cpp new file mode 100644 index 00000000000..0abceaa14df --- /dev/null +++ b/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 +#include + +#include "json_object.h" + +namespace cldnn { + +primitive_type_id experimental_detectron_prior_grid_generator::type_id() { + static primitive_type_base 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(desc->featmap_width * desc->featmap_height * data_layout.batch()), 4, 1, 1}); + } else { + return layout(data_layout.data_type, + format::bfyx, + {static_cast(desc->featmap_height), + static_cast(desc->featmap_width), + 4, + static_cast(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 diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_prior_grid_generator.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_prior_grid_generator.cpp new file mode 100644 index 00000000000..a0b3d373aa9 --- /dev/null +++ b/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 +#include +#include +#include + +#include +#include + +#include "primitive_base.hpp" + +namespace cldnn { +namespace ocl { + +struct experimental_detectron_prior_grid_generator_impl + : typed_primitive_impl_ocl { + using parent = typed_primitive_impl_ocl; + using parent::parent; + + std::unique_ptr clone() const override { + return make_unique(*this); + } + + static primitive_impl* create(const experimental_detectron_prior_grid_generator_node& arg) { + auto params = get_default_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(prim.image_width) / params.layer_width; + params.step_y = prim.stride_y ? prim.stride_y : static_cast(prim.image_height) / params.layer_height; + + auto optional_params = + get_default_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::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 diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/register.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/register.cpp index 2f217052d06..e93fddc996d 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/register.cpp +++ b/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); diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/register.hpp b/src/plugins/intel_gpu/src/graph/impls/ocl/register.hpp index 9c391b9c377..545da43e52c 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/register.hpp +++ b/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); diff --git a/src/plugins/intel_gpu/src/graph/include/experimental_detectron_prior_grid_generator_inst.h b/src/plugins/intel_gpu/src/graph/include/experimental_detectron_prior_grid_generator_inst.h new file mode 100644 index 00000000000..b4d72799fb9 --- /dev/null +++ b/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 + +#include "primitive_inst.h" + +namespace cldnn { +template <> +struct typed_program_node + : public typed_program_node_base { + 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; + +template <> +class typed_primitive_inst + : public typed_primitive_inst_base { +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; + +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/kernel_selector/common/common_types.h b/src/plugins/intel_gpu/src/kernel_selector/common/common_types.h index 6d7d0f9a988..a834b329379 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/common/common_types.h +++ b/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, diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/edpgg/experimental_detectron_prior_grid_generator_kernel_ref.cpp b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/edpgg/experimental_detectron_prior_grid_generator_kernel_ref.cpp new file mode 100644 index 00000000000..186594decc9 --- /dev/null +++ b/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 + +namespace kernel_selector { + +namespace { + +CommonDispatchData SetDefault(const experimental_detectron_prior_grid_generator_params ¶ms) { + 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 ¶ms, const optional_params &options) const { + KernelsData kernels_data; + if (!Validate(params, options)) + return kernels_data; + kernels_data.push_back(KernelData::Default(params)); + KernelData &kernel_data = kernels_data.front(); + auto &derived_params = dynamic_cast(*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 ¶ms = dynamic_cast(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 diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/edpgg/experimental_detectron_prior_grid_generator_kernel_ref.h b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/edpgg/experimental_detectron_prior_grid_generator_kernel_ref.h new file mode 100644 index 00000000000..279da170e5c --- /dev/null +++ b/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 ¶ms, const optional_params &options) const override; + + KernelsPriority GetKernelsPriority(const Params ¶ms, 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 ¶ms) const; +}; + +} // namespace kernel_selector diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/edpgg/experimental_detectron_prior_grid_generator_kernel_selector.cpp b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/edpgg/experimental_detectron_prior_grid_generator_kernel_selector.cpp new file mode 100644 index 00000000000..0403821853d --- /dev/null +++ b/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 ¶ms, 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(); + } +}; + +} // namespace + +kernel_selector_base& experimental_detectron_prior_grid_generator_instance() { + static experimental_detectron_prior_grid_generator_kernel_selector instance; + return instance; +} + +} // namespace kernel_selector diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/edpgg/experimental_detectron_prior_grid_generator_kernel_selector.h b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/edpgg/experimental_detectron_prior_grid_generator_kernel_selector.h new file mode 100644 index 00000000000..e845872cfb5 --- /dev/null +++ b/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 + +namespace kernel_selector { + +/* + * ExperimentalDetectronPriorGridGenerator kernel selector + */ +kernel_selector_base& experimental_detectron_prior_grid_generator_instance(); + +} // namespace kernel_selector diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/experimental_detectron_prior_grid_generator_ref.cl b/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/experimental_detectron_prior_grid_generator_ref.cl new file mode 100644 index 00000000000..ff1264b7467 --- /dev/null +++ b/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 diff --git a/src/plugins/intel_gpu/src/plugin/ops/experimental_detectron_prior_grid_generator.cpp b/src/plugins/intel_gpu/src/plugin/ops/experimental_detectron_prior_grid_generator.cpp new file mode 100644 index 00000000000..d82d22829a5 --- /dev/null +++ b/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 +#include +#include +#include + +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& 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(attrs.h), + static_cast(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 diff --git a/src/plugins/intel_gpu/tests/test_cases/experimental_detectron_prior_grid_generator_gpu_test.cpp b/src/plugins/intel_gpu/tests/test_cases/experimental_detectron_prior_grid_generator_gpu_test.cpp new file mode 100644 index 00000000000..0f6827533e7 --- /dev/null +++ b/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 +#include +#include +#include +#include + +#include "test_utils.h" + +using namespace cldnn; +using namespace ::tests; + +template +struct ExperimentalDetectronPriorGridGeneratorParams { + std::vector priors; + tensor priorsTensor; + int h; + int w; + float strideX; + float strideY; + bool flatten; + std::pair featureShape; + std::pair imageShape; + tensor outputTensor; + std::vector expectedOutput; +}; + +template +struct experimental_detectron_prior_grid_generator_test + : public ::testing::TestWithParam> { +public: + void test() { + auto data_type = type_to_data_type::value; + ExperimentalDetectronPriorGridGeneratorParams params = + testing::TestWithParam>::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 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 +std::vector getValues(const std::vector& values) { + std::vector result(values.begin(), values.end()); + return result; +} + +template +std::vector> generateExperimentalPGGParams() { + std::vector> experimentalPGGParams{ + {getValues({-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( + {-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({-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( + {-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({-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({-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({-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({-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 + std::string operator()(const testing::TestParamInfo>& 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; +using experimental_detectron_prior_grid_generator_test_f16 = experimental_detectron_prior_grid_generator_test; + +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()), + PrintToStringParamName()); + +INSTANTIATE_TEST_SUITE_P(smoke_experimental_detectron_prior_grid_generator_test_f16, + experimental_detectron_prior_grid_generator_test_f16, + ::testing::ValuesIn(generateExperimentalPGGParams()), + PrintToStringParamName()); diff --git a/src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/experimental_detectron_prior_grid_generator.cpp b/src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/experimental_detectron_prior_grid_generator.cpp new file mode 100644 index 00000000000..d644dbe7326 --- /dev/null +++ b/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 +#include +#include + +namespace { + +const std::initializer_list 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 +std::vector>> getInputTensors() { + std::vector>> tensors{ + {"test#1", + {ov::test::utils::create_tensor( + ov::element::from(), + {3, 4}, + std::vector{-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( + ov::element::from(), + {3, 4}, + std::vector{-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( + ov::element::from(), + {3, 4}, + std:: + vector{-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( + ov::element::from(), + {3, 4}, + std::vector{-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()), + 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()), + testing::ValuesIn({ov::element::Type_t::f16}), + testing::Values(CommonTestUtils::DEVICE_GPU)), + ExperimentalDetectronPriorGridGeneratorLayerTest::getTestCaseName); +} // namespace diff --git a/src/tests/functional/shared_test_classes/src/single_layer/experimental_detectron_prior_grid_generator.cpp b/src/tests/functional/shared_test_classes/src/single_layer/experimental_detectron_prior_grid_generator.cpp index 6f8fb0e4225..d8c8142da17 100644 --- a/src/tests/functional/shared_test_classes/src/single_layer/experimental_detectron_prior_grid_generator.cpp +++ b/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 +ov::runtime::Tensor generateTensorByShape(const Shape &shape) { + return ov::test::utils::create_tensor( + ov::element::from(), + shape, + std::vector(0., shape_size(shape))); +} +} + void ExperimentalDetectronPriorGridGeneratorLayerTest::generate_inputs(const std::vector& 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( - ov::element::f32, - targetInputStaticShapes[j], - std::vector(0.f, shape_size(targetInputStaticShapes[j]))); + ov::runtime::Tensor inputTensor = (netPrecision == element::f16) + ? generateTensorByShape(targetInputStaticShapes[j]) + : generateTensorByShape( + targetInputStaticShapes[j]); inputs.insert({funcInputs[j].get_node_shared_ptr(), inputTensor}); }