From 3de00347f3a6ae75ccaa4e935e50f8002dac026d Mon Sep 17 00:00:00 2001 From: Vladimir Paramuzov Date: Wed, 1 Mar 2023 14:05:59 +0400 Subject: [PATCH] [GPU] Code cleanup (#16014) * [GPU] Improve exception message for program build * [GPU] Code cleanup --- .../intel_gpu/src/graph/adaptive_pooling.cpp | 1 - .../intel_gpu/src/graph/arg_max_min.cpp | 1 - src/plugins/intel_gpu/src/graph/assign.cpp | 3 +- .../intel_gpu/src/graph/batch_to_space.cpp | 1 - .../src/graph/binary_convolution.cpp | 1 - .../intel_gpu/src/graph/convert_color.cpp | 1 - src/plugins/intel_gpu/src/graph/ctc_loss.cpp | 1 - src/plugins/intel_gpu/src/graph/cum_sum.cpp | 1 - .../src/graph/deformable_convolution.cpp | 1 - ...xperimental_detectron_detection_output.cpp | 1 - ...tectron_generate_proposal_single_image.cpp | 1 - ...mental_detectron_roi_feature_extractor.cpp | 1 - .../experimental_detectron_topk_rois.cpp | 1 - .../src/graph/extract_image_patches.cpp | 1 - .../intel_gpu/src/graph/fully_connected.cpp | 1 - src/plugins/intel_gpu/src/graph/gather.cpp | 1 - .../intel_gpu/src/graph/gather_elements.cpp | 1 - src/plugins/intel_gpu/src/graph/gather_nd.cpp | 1 - src/plugins/intel_gpu/src/graph/gemm.cpp | 1 - .../src/graph/generate_proposals.cpp | 1 - .../graph_optimizer/add_required_reorders.cpp | 2 + .../graph_optimizer/graph_initializations.cpp | 1 - .../graph_optimizer/post_input_reorder.cpp | 2 + .../graph_optimizer/pre_replace_deconv.cpp | 4 +- .../prepare_primitive_fusing.cpp | 3 +- .../graph_optimizer/prepare_quantization.cpp | 2 + .../src/graph/impls/common/condition.cpp | 2 +- .../intel_gpu/src/graph/impls/common/loop.cpp | 3 +- .../graph/impls/common/wait_for_events.cpp | 4 +- .../intel_gpu/src/graph/impls/cpu/assign.cpp | 3 +- .../src/graph/impls/cpu/detection_output.cpp | 21 +- .../graph/impls/cpu/non_max_suppression.cpp | 2 +- .../src/graph/impls/cpu/proposal.cpp | 2 +- .../src/graph/impls/cpu/read_value.cpp | 3 +- .../src/graph/impls/ocl/activation.cpp | 12 +- .../src/graph/impls/ocl/adaptive_pooling.cpp | 7 +- .../src/graph/impls/ocl/arg_max_min.cpp | 6 +- .../src/graph/impls/ocl/batch_to_space.cpp | 9 +- .../graph/impls/ocl/binary_convolution.cpp | 9 +- .../intel_gpu/src/graph/impls/ocl/border.cpp | 6 +- .../src/graph/impls/ocl/broadcast.cpp | 7 +- .../src/graph/impls/ocl/bucketize.cpp | 6 +- .../src/graph/impls/ocl/concatenation.cpp | 8 +- .../src/graph/impls/ocl/convert_color.cpp | 10 +- .../src/graph/impls/ocl/convolution.cpp | 9 +- .../intel_gpu/src/graph/impls/ocl/crop.cpp | 6 +- .../graph/impls/ocl/ctc_greedy_decoder.cpp | 10 +- .../src/graph/impls/ocl/ctc_loss.cpp | 6 +- .../intel_gpu/src/graph/impls/ocl/cum_sum.cpp | 7 +- .../src/graph/impls/ocl/custom_primitive.cpp | 11 +- .../src/graph/impls/ocl/deconvolution.cpp | 7 +- .../impls/ocl/deformable_convolution.cpp | 7 +- .../src/graph/impls/ocl/depth_to_space.cpp | 9 +- .../src/graph/impls/ocl/detection_output.cpp | 7 +- .../intel_gpu/src/graph/impls/ocl/dft.cpp | 12 +- .../intel_gpu/src/graph/impls/ocl/eltwise.cpp | 7 +- .../src/graph/impls/ocl/embedding_bag.cpp | 9 +- ...xperimental_detectron_detection_output.cpp | 7 +- ...ectron_generate_proposals_single_image.cpp | 6 +- ...imental_detectron_prior_grid_generator.cpp | 14 +- ...mental_detectron_roi_feature_extractor.cpp | 6 +- .../ocl/experimental_detectron_topk_rois.cpp | 10 +- .../graph/impls/ocl/extract_image_patches.cpp | 5 +- .../intel_gpu/src/graph/impls/ocl/eye.cpp | 15 +- .../src/graph/impls/ocl/fully_connected.cpp | 12 +- .../intel_gpu/src/graph/impls/ocl/gather.cpp | 8 +- .../src/graph/impls/ocl/gather_elements.cpp | 8 +- .../src/graph/impls/ocl/gather_nd.cpp | 7 +- .../src/graph/impls/ocl/gather_tree.cpp | 7 +- .../intel_gpu/src/graph/impls/ocl/gemm.cpp | 9 +- .../graph/impls/ocl/generate_proposals.cpp | 7 +- .../src/graph/impls/ocl/generic_layer.cpp | 7 +- .../src/graph/impls/ocl/grid_sample.cpp | 6 +- .../intel_gpu/src/graph/impls/ocl/grn.cpp | 10 +- .../ocl}/kernel_selector_helper.cpp | 7 + .../ocl}/kernel_selector_helper.h | 120 +----- .../intel_gpu/src/graph/impls/ocl/lrn.cpp | 6 +- .../graph/impls/ocl/lstm_dynamic_input.cpp | 6 +- .../graph/impls/ocl/lstm_dynamic_timeloop.cpp | 6 +- .../src/graph/impls/ocl/lstm_elt.cpp | 61 ++- .../src/graph/impls/ocl/lstm_gemm.cpp | 6 +- .../src/graph/impls/ocl/matrix_nms.cpp | 8 +- .../src/graph/impls/ocl/multiclass_nms.cpp | 6 +- .../src/graph/impls/ocl/mutable_data.cpp | 4 +- .../intel_gpu/src/graph/impls/ocl/mvn.cpp | 10 +- .../graph/impls/ocl/non_max_suppression.cpp | 7 +- .../src/graph/impls/ocl/non_zero.cpp | 8 +- .../src/graph/impls/ocl/normalize.cpp | 10 +- .../intel_gpu/src/graph/impls/ocl/one_hot.cpp | 8 +- .../intel_gpu/src/graph/impls/ocl/permute.cpp | 8 +- .../intel_gpu/src/graph/impls/ocl/pooling.cpp | 7 +- .../src/graph/impls/ocl/primitive_base.hpp | 11 +- .../src/graph/impls/ocl/prior_box.cpp | 12 +- .../src/graph/impls/ocl/pyramid_roi_align.cpp | 6 +- .../src/graph/impls/ocl/quantize.cpp | 8 +- .../src/graph/impls/ocl/random_uniform.cpp | 10 +- .../intel_gpu/src/graph/impls/ocl/range.cpp | 10 +- .../intel_gpu/src/graph/impls/ocl/reduce.cpp | 10 +- .../src/graph/impls/ocl/region_yolo.cpp | 6 +- .../intel_gpu/src/graph/impls/ocl/reorder.cpp | 6 +- .../src/graph/impls/ocl/reorg_yolo.cpp | 6 +- .../src/graph/impls/ocl/resample.cpp | 8 +- .../intel_gpu/src/graph/impls/ocl/reshape.cpp | 6 +- .../intel_gpu/src/graph/impls/ocl/reverse.cpp | 10 +- .../src/graph/impls/ocl/reverse_sequence.cpp | 8 +- .../src/graph/impls/ocl/roi_align.cpp | 6 +- .../src/graph/impls/ocl/roi_pooling.cpp | 6 +- .../intel_gpu/src/graph/impls/ocl/roll.cpp | 7 +- .../impls/ocl/scatter_elements_update.cpp | 8 +- .../src/graph/impls/ocl/scatter_nd_update.cpp | 8 +- .../src/graph/impls/ocl/scatter_update.cpp | 8 +- .../intel_gpu/src/graph/impls/ocl/select.cpp | 6 +- .../src/graph/impls/ocl/shape_of.cpp | 6 +- .../src/graph/impls/ocl/shuffle_channels.cpp | 8 +- .../intel_gpu/src/graph/impls/ocl/slice.cpp | 20 +- .../intel_gpu/src/graph/impls/ocl/softmax.cpp | 6 +- .../src/graph/impls/ocl/space_to_batch.cpp | 10 +- .../src/graph/impls/ocl/space_to_depth.cpp | 8 +- .../src/graph/impls/ocl/strided_slice.cpp | 11 +- .../intel_gpu/src/graph/impls/ocl/tile.cpp | 8 +- .../impls/onednn/concatenation_onednn.cpp | 2 +- .../graph/impls/onednn/convolution_onednn.cpp | 2 +- .../impls/onednn/deconvolution_onednn.cpp | 2 +- .../src/graph/impls/onednn/eltwise_onednn.cpp | 2 +- .../impls/onednn/fully_connected_onednn.cpp | 2 +- .../src/graph/impls/onednn/gemm_onednn.cpp | 2 +- .../src/graph/impls/onednn/pooling_onednn.cpp | 2 +- .../impls/onednn/primitive_onednn_base.h | 1 - .../graph/impls/onednn/reduction_onednn.cpp | 2 +- .../src/graph/impls/onednn/reorder_onednn.cpp | 2 +- .../intel_gpu/src/graph/include/assign_inst.h | 1 - .../experimental_detectron_topk_rois_inst.h | 1 - .../intel_gpu/src/graph/include/eye_inst.h | 1 - .../src/graph/include/generic_layer.hpp | 6 +- .../{impls => include}/implementation_map.hpp | 31 +- .../src/graph/include/input_layout_inst.h | 1 - .../src/graph/include/kernel_impl_params.hpp | 124 ++++++ .../src/graph/include/layout_optimizer.h | 4 +- .../intel_gpu/src/graph/include/loop_inst.h | 20 +- .../graph/include/lstm_dynamic_input_inst.h | 1 + .../src/graph/include/lstm_dynamic_inst.h | 1 + .../src/graph/include/non_zero_inst.h | 1 - .../src/graph/include/primitive_inst.h | 7 +- .../src/graph/include/primitive_type.h | 2 +- .../src/graph/include/primitive_type_base.h | 22 +- .../src/graph/include/program_node.h | 7 +- .../src/graph/include/random_uniform_inst.h | 1 - .../src/graph/include/read_value_inst.h | 1 - .../src/graph/include/reshape_inst.h | 6 +- .../src/graph/include/roi_align_inst.h | 1 - .../intel_gpu/src/graph/include/slice_inst.h | 1 - .../graph/include/sliding_window_utils.hpp | 2 +- .../include/sliding_window_utils_legacy.h | 389 ------------------ .../src/graph/include/strided_slice_inst.h | 1 - .../src/graph/include/to_string_utils.h | 16 - .../intel_gpu/src/graph/input_layout.cpp | 2 +- src/plugins/intel_gpu/src/graph/loop.cpp | 1 + src/plugins/intel_gpu/src/graph/lstm_gemm.cpp | 2 +- .../intel_gpu/src/graph/matrix_nms.cpp | 2 + .../intel_gpu/src/graph/multiclass_nms.cpp | 1 - .../intel_gpu/src/graph/mutable_data.cpp | 2 - src/plugins/intel_gpu/src/graph/network.cpp | 14 +- .../intel_gpu/src/graph/nodes_ordering.cpp | 1 - src/plugins/intel_gpu/src/graph/non_zero.cpp | 2 +- .../intel_gpu/src/graph/pass_manager.cpp | 2 + .../intel_gpu/src/graph/primitive_inst.cpp | 2 +- src/plugins/intel_gpu/src/graph/program.cpp | 9 +- .../src/graph/program_dump_graph.cpp | 2 +- .../intel_gpu/src/graph/program_node.cpp | 2 +- .../intel_gpu/src/graph/pyramid_roi_align.cpp | 1 - src/plugins/intel_gpu/src/graph/quantize.cpp | 1 - .../intel_gpu/src/graph/read_value.cpp | 1 - src/plugins/intel_gpu/src/graph/reduce.cpp | 2 - src/plugins/intel_gpu/src/graph/resample.cpp | 1 - .../intel_gpu/src/graph/reverse_sequence.cpp | 1 - src/plugins/intel_gpu/src/graph/roi_align.cpp | 2 + .../intel_gpu/src/graph/roi_pooling.cpp | 1 - .../src/graph/scatter_elements_update.cpp | 2 +- .../intel_gpu/src/graph/scatter_nd_update.cpp | 1 - .../intel_gpu/src/graph/scatter_update.cpp | 1 - src/plugins/intel_gpu/src/graph/shape_of.cpp | 1 - .../intel_gpu/src/graph/space_to_depth.cpp | 2 +- .../intel_gpu/src/graph/strided_slice.cpp | 2 - src/plugins/intel_gpu/src/graph/tile.cpp | 1 - src/plugins/intel_gpu/src/plugin/program.cpp | 2 +- .../ocl/ocl_command_queues_builder.cpp | 1 - .../intel_gpu/tests/test_utils/network_test.h | 1 + .../intel_gpu/tests/test_utils/test_utils.h | 38 +- ...compose_reduce_for_false_keepdims_test.cpp | 1 + 189 files changed, 501 insertions(+), 1211 deletions(-) rename src/plugins/intel_gpu/src/graph/{ => impls/ocl}/kernel_selector_helper.cpp (99%) rename src/plugins/intel_gpu/src/graph/{include => impls/ocl}/kernel_selector_helper.h (72%) rename src/plugins/intel_gpu/src/graph/{impls => include}/implementation_map.hpp (91%) create mode 100644 src/plugins/intel_gpu/src/graph/include/kernel_impl_params.hpp delete mode 100644 src/plugins/intel_gpu/src/graph/include/sliding_window_utils_legacy.h diff --git a/src/plugins/intel_gpu/src/graph/adaptive_pooling.cpp b/src/plugins/intel_gpu/src/graph/adaptive_pooling.cpp index bf1a39d5710..26778bc232f 100644 --- a/src/plugins/intel_gpu/src/graph/adaptive_pooling.cpp +++ b/src/plugins/intel_gpu/src/graph/adaptive_pooling.cpp @@ -4,7 +4,6 @@ #include "adaptive_pooling_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/arg_max_min.cpp b/src/plugins/intel_gpu/src/graph/arg_max_min.cpp index ca79e4dee67..21665eceafc 100644 --- a/src/plugins/intel_gpu/src/graph/arg_max_min.cpp +++ b/src/plugins/intel_gpu/src/graph/arg_max_min.cpp @@ -4,7 +4,6 @@ #include "arg_max_min_inst.h" #include "primitive_type_base.h" -#include "sliding_window_utils_legacy.h" #include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/assign.cpp b/src/plugins/intel_gpu/src/graph/assign.cpp index 8ef48a7e2f7..b9f57e7c4d1 100644 --- a/src/plugins/intel_gpu/src/graph/assign.cpp +++ b/src/plugins/intel_gpu/src/graph/assign.cpp @@ -2,11 +2,10 @@ // SPDX-License-Identifier: Apache-2.0 // -#include +#include "assign_inst.h" #include "primitive_type_base.h" #include #include -#include namespace cldnn { GPU_DEFINE_PRIMITIVE_TYPE_ID(assign) diff --git a/src/plugins/intel_gpu/src/graph/batch_to_space.cpp b/src/plugins/intel_gpu/src/graph/batch_to_space.cpp index 9166dc58be9..c6adbd34802 100644 --- a/src/plugins/intel_gpu/src/graph/batch_to_space.cpp +++ b/src/plugins/intel_gpu/src/graph/batch_to_space.cpp @@ -7,7 +7,6 @@ #include "primitive_type_base.h" #include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" -#include "data_inst.h" #include #include diff --git a/src/plugins/intel_gpu/src/graph/binary_convolution.cpp b/src/plugins/intel_gpu/src/graph/binary_convolution.cpp index 2c49f04dc4b..ff8a34c3364 100644 --- a/src/plugins/intel_gpu/src/graph/binary_convolution.cpp +++ b/src/plugins/intel_gpu/src/graph/binary_convolution.cpp @@ -3,7 +3,6 @@ // #include "binary_convolution_inst.h" #include "convolution_inst.h" -#include "reorder_inst.h" #include "primitive_type_base.h" #include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" diff --git a/src/plugins/intel_gpu/src/graph/convert_color.cpp b/src/plugins/intel_gpu/src/graph/convert_color.cpp index 9db657884be..8e9bc54dd58 100644 --- a/src/plugins/intel_gpu/src/graph/convert_color.cpp +++ b/src/plugins/intel_gpu/src/graph/convert_color.cpp @@ -5,7 +5,6 @@ #include "convert_color_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/ctc_loss.cpp b/src/plugins/intel_gpu/src/graph/ctc_loss.cpp index 52d3bdbdc38..3a57a1caf2f 100644 --- a/src/plugins/intel_gpu/src/graph/ctc_loss.cpp +++ b/src/plugins/intel_gpu/src/graph/ctc_loss.cpp @@ -8,7 +8,6 @@ #include "ctc_loss_inst.hpp" #include "json_object.h" #include "primitive_type_base.h" -#include "to_string_utils.h" namespace cldnn { GPU_DEFINE_PRIMITIVE_TYPE_ID(ctc_loss) diff --git a/src/plugins/intel_gpu/src/graph/cum_sum.cpp b/src/plugins/intel_gpu/src/graph/cum_sum.cpp index 926207c6083..4a4d4c022d7 100644 --- a/src/plugins/intel_gpu/src/graph/cum_sum.cpp +++ b/src/plugins/intel_gpu/src/graph/cum_sum.cpp @@ -5,7 +5,6 @@ #include "cum_sum_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/deformable_convolution.cpp b/src/plugins/intel_gpu/src/graph/deformable_convolution.cpp index ff25917cc2b..2c7cf6d5425 100644 --- a/src/plugins/intel_gpu/src/graph/deformable_convolution.cpp +++ b/src/plugins/intel_gpu/src/graph/deformable_convolution.cpp @@ -3,7 +3,6 @@ // #include "deformable_convolution_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/experimental_detectron_detection_output.cpp b/src/plugins/intel_gpu/src/graph/experimental_detectron_detection_output.cpp index 63cf986939a..d375002cb2f 100644 --- a/src/plugins/intel_gpu/src/graph/experimental_detectron_detection_output.cpp +++ b/src/plugins/intel_gpu/src/graph/experimental_detectron_detection_output.cpp @@ -5,7 +5,6 @@ #include #include "experimental_detectron_detection_output_inst.hpp" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include "primitive_type_base.h" diff --git a/src/plugins/intel_gpu/src/graph/experimental_detectron_generate_proposal_single_image.cpp b/src/plugins/intel_gpu/src/graph/experimental_detectron_generate_proposal_single_image.cpp index 4399a9b4d12..413d2fe0400 100644 --- a/src/plugins/intel_gpu/src/graph/experimental_detectron_generate_proposal_single_image.cpp +++ b/src/plugins/intel_gpu/src/graph/experimental_detectron_generate_proposal_single_image.cpp @@ -4,7 +4,6 @@ #include "experimental_detectron_generate_proposals_single_image_inst.hpp" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/experimental_detectron_roi_feature_extractor.cpp b/src/plugins/intel_gpu/src/graph/experimental_detectron_roi_feature_extractor.cpp index 804c42ab052..9bdbf5747c1 100644 --- a/src/plugins/intel_gpu/src/graph/experimental_detectron_roi_feature_extractor.cpp +++ b/src/plugins/intel_gpu/src/graph/experimental_detectron_roi_feature_extractor.cpp @@ -4,7 +4,6 @@ #include "experimental_detectron_roi_feature_extractor_inst.hpp" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/experimental_detectron_topk_rois.cpp b/src/plugins/intel_gpu/src/graph/experimental_detectron_topk_rois.cpp index 2eced6861a4..d08e181e9e3 100644 --- a/src/plugins/intel_gpu/src/graph/experimental_detectron_topk_rois.cpp +++ b/src/plugins/intel_gpu/src/graph/experimental_detectron_topk_rois.cpp @@ -6,7 +6,6 @@ #include "primitive_type_base.h" #include #include -#include namespace cldnn { GPU_DEFINE_PRIMITIVE_TYPE_ID(experimental_detectron_topk_rois) diff --git a/src/plugins/intel_gpu/src/graph/extract_image_patches.cpp b/src/plugins/intel_gpu/src/graph/extract_image_patches.cpp index c9b1861a5e4..273a2f7501c 100644 --- a/src/plugins/intel_gpu/src/graph/extract_image_patches.cpp +++ b/src/plugins/intel_gpu/src/graph/extract_image_patches.cpp @@ -5,7 +5,6 @@ #include "extract_image_patches_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/fully_connected.cpp b/src/plugins/intel_gpu/src/graph/fully_connected.cpp index d00a36c676b..80a3b97b844 100644 --- a/src/plugins/intel_gpu/src/graph/fully_connected.cpp +++ b/src/plugins/intel_gpu/src/graph/fully_connected.cpp @@ -3,7 +3,6 @@ // #include "fully_connected_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include #include diff --git a/src/plugins/intel_gpu/src/graph/gather.cpp b/src/plugins/intel_gpu/src/graph/gather.cpp index cf68b5e98d7..7c022b1be2b 100644 --- a/src/plugins/intel_gpu/src/graph/gather.cpp +++ b/src/plugins/intel_gpu/src/graph/gather.cpp @@ -5,7 +5,6 @@ #include "gather_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/gather_elements.cpp b/src/plugins/intel_gpu/src/graph/gather_elements.cpp index ef1b0cc7c38..af91f9fdf4d 100644 --- a/src/plugins/intel_gpu/src/graph/gather_elements.cpp +++ b/src/plugins/intel_gpu/src/graph/gather_elements.cpp @@ -6,7 +6,6 @@ #include "gather_elements_shape_inference.hpp" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/gather_nd.cpp b/src/plugins/intel_gpu/src/graph/gather_nd.cpp index f2d00eea38e..e5295573e40 100644 --- a/src/plugins/intel_gpu/src/graph/gather_nd.cpp +++ b/src/plugins/intel_gpu/src/graph/gather_nd.cpp @@ -5,7 +5,6 @@ #include "gather_nd_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/gemm.cpp b/src/plugins/intel_gpu/src/graph/gemm.cpp index b15be72fc4d..a43148c7ed9 100644 --- a/src/plugins/intel_gpu/src/graph/gemm.cpp +++ b/src/plugins/intel_gpu/src/graph/gemm.cpp @@ -3,7 +3,6 @@ // #include "gemm_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include #include diff --git a/src/plugins/intel_gpu/src/graph/generate_proposals.cpp b/src/plugins/intel_gpu/src/graph/generate_proposals.cpp index f0a13313f3f..5822a1fef09 100644 --- a/src/plugins/intel_gpu/src/graph/generate_proposals.cpp +++ b/src/plugins/intel_gpu/src/graph/generate_proposals.cpp @@ -4,7 +4,6 @@ #include "generate_proposals_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/graph_optimizer/add_required_reorders.cpp b/src/plugins/intel_gpu/src/graph/graph_optimizer/add_required_reorders.cpp index 439f144f8f0..9c866ba343f 100644 --- a/src/plugins/intel_gpu/src/graph/graph_optimizer/add_required_reorders.cpp +++ b/src/plugins/intel_gpu/src/graph/graph_optimizer/add_required_reorders.cpp @@ -13,6 +13,8 @@ #include #include +using namespace cldnn; + /* This pass checks if data formats (layouts) of output/input in hidden layers match. If not than required reorder is added to the network. diff --git a/src/plugins/intel_gpu/src/graph/graph_optimizer/graph_initializations.cpp b/src/plugins/intel_gpu/src/graph/graph_optimizer/graph_initializations.cpp index df9940e8110..c69b05710ec 100644 --- a/src/plugins/intel_gpu/src/graph/graph_optimizer/graph_initializations.cpp +++ b/src/plugins/intel_gpu/src/graph/graph_optimizer/graph_initializations.cpp @@ -17,7 +17,6 @@ #include "lstm_dynamic_timeloop_inst.h" #include "mutable_data_inst.h" #include "arg_max_min_inst.h" -#include "kernel_selector_utils.h" #include #include diff --git a/src/plugins/intel_gpu/src/graph/graph_optimizer/post_input_reorder.cpp b/src/plugins/intel_gpu/src/graph/graph_optimizer/post_input_reorder.cpp index 8a3c93f9007..49042549132 100644 --- a/src/plugins/intel_gpu/src/graph/graph_optimizer/post_input_reorder.cpp +++ b/src/plugins/intel_gpu/src/graph/graph_optimizer/post_input_reorder.cpp @@ -8,6 +8,8 @@ #include #include +using namespace cldnn; + /* This pass checks if if primitive's input format matches implementation's input format If not than required reorder is added to the network. diff --git a/src/plugins/intel_gpu/src/graph/graph_optimizer/pre_replace_deconv.cpp b/src/plugins/intel_gpu/src/graph/graph_optimizer/pre_replace_deconv.cpp index 2fa294fe8be..d68437bf406 100644 --- a/src/plugins/intel_gpu/src/graph/graph_optimizer/pre_replace_deconv.cpp +++ b/src/plugins/intel_gpu/src/graph/graph_optimizer/pre_replace_deconv.cpp @@ -8,13 +8,13 @@ #include "convolution_inst.h" #include "deconvolution_inst.h" #include "depth_to_space_inst.h" -#include "kernel_selector_utils.h" #include #include #include #include #include -#include "intel_gpu/runtime/error_handler.hpp" + +using namespace cldnn; void pre_replace_deconv::run(program& p) { bool update_processing_order = false; diff --git a/src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_primitive_fusing.cpp b/src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_primitive_fusing.cpp index bfe7d97fe28..8f68450bf49 100644 --- a/src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_primitive_fusing.cpp +++ b/src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_primitive_fusing.cpp @@ -47,11 +47,12 @@ #include #include #include -#include "intel_gpu/runtime/error_handler.hpp" #ifdef ENABLE_ONEDNN_FOR_GPU #include #endif +using namespace cldnn; + void prepare_primitive_fusing::run(program& p) { fuse_reorders(p); remove_redundant_reshape(p); diff --git a/src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_quantization.cpp b/src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_quantization.cpp index 167d58679a1..7a290c19c67 100644 --- a/src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_quantization.cpp +++ b/src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_quantization.cpp @@ -19,6 +19,8 @@ #include #include +using namespace cldnn; + template bool check_binarization(memory::ptr mem_input_low, memory::ptr mem_input_high, program& p) { bool is_binarization = true; diff --git a/src/plugins/intel_gpu/src/graph/impls/common/condition.cpp b/src/plugins/intel_gpu/src/graph/impls/common/condition.cpp index f5417c79138..c760faa8a43 100644 --- a/src/plugins/intel_gpu/src/graph/impls/common/condition.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/common/condition.cpp @@ -3,7 +3,7 @@ // #include "condition_inst.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "register.hpp" #include diff --git a/src/plugins/intel_gpu/src/graph/impls/common/loop.cpp b/src/plugins/intel_gpu/src/graph/impls/common/loop.cpp index 27ea5e25775..c52cd78e2b1 100644 --- a/src/plugins/intel_gpu/src/graph/impls/common/loop.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/common/loop.cpp @@ -2,11 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // #include "loop_inst.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "register.hpp" #include "mutable_data_inst.h" #include "input_layout_inst.h" #include "intel_gpu/graph/serialization/loop_serializer.hpp" +#include "intel_gpu/runtime/error_handler.hpp" #include #include diff --git a/src/plugins/intel_gpu/src/graph/impls/common/wait_for_events.cpp b/src/plugins/intel_gpu/src/graph/impls/common/wait_for_events.cpp index ef63f3c2fb5..2a9674bb8ad 100644 --- a/src/plugins/intel_gpu/src/graph/impls/common/wait_for_events.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/common/wait_for_events.cpp @@ -6,7 +6,7 @@ #include "data_inst.h" #include "prior_box_inst.h" #include "input_layout_inst.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "register.hpp" #include "intel_gpu/graph/serialization/binary_buffer.hpp" #include @@ -19,7 +19,7 @@ class wait_for_events_impl : public primitive_impl { public: explicit wait_for_events_impl(const program_node& /*node*/) - : primitive_impl(kernel_selector::weights_reorder_params{}, "wait_for_events") { } + : primitive_impl("wait_for_events") { } wait_for_events_impl() : primitive_impl() {} diff --git a/src/plugins/intel_gpu/src/graph/impls/cpu/assign.cpp b/src/plugins/intel_gpu/src/graph/impls/cpu/assign.cpp index 07bbef6816b..fea869d1ac3 100644 --- a/src/plugins/intel_gpu/src/graph/impls/cpu/assign.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/cpu/assign.cpp @@ -3,8 +3,9 @@ // #include "assign_inst.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "register.hpp" +#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace cpu { diff --git a/src/plugins/intel_gpu/src/graph/impls/cpu/detection_output.cpp b/src/plugins/intel_gpu/src/graph/impls/cpu/detection_output.cpp index 93fad828c55..14a3e15b1a1 100644 --- a/src/plugins/intel_gpu/src/graph/impls/cpu/detection_output.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/cpu/detection_output.cpp @@ -3,7 +3,7 @@ // #include "detection_output_inst.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "register.hpp" #include "cpu_impl_helpers.hpp" @@ -16,19 +16,9 @@ #include #include -#ifdef FIX_OPENMP_RELEASE_ISSUE -#ifdef OPENMP_FOUND -#include -#endif -#endif - namespace cldnn { namespace cpu { -namespace { - using bounding_box = cldnn::cpu::bounding_box; -} // namespace - template bool comp_score_descend(const std::pair& pair1, const std::pair& pair2) { @@ -296,15 +286,6 @@ public: std::vector>>& conf_per_image = confidences[image]; std::map> indices; int num_det = 0; -#ifdef FIX_OPENMP_RELEASE_ISSUE -#ifdef OPENMP_FOUND - int num_available_threads = omp_get_max_threads(); - // half available threads usage shows the best perf results for both SKL (4c8t) and APL (4c4t) for this part - // of detection output - int num_threads_to_use = (omp_in_parallel() == 0) ? num_available_threads / 2 : 1; -#pragma omp parallel for num_threads(num_threads_to_use) reduction(+ : num_det) -#endif -#endif if (nms_type == NMSType::CAFFE) { for (int cls = 0; cls < static_cast(args->num_classes); ++cls) { if (static_cast(cls) == args->background_label_id) { diff --git a/src/plugins/intel_gpu/src/graph/impls/cpu/non_max_suppression.cpp b/src/plugins/intel_gpu/src/graph/impls/cpu/non_max_suppression.cpp index 2a6bab53ec6..c5c87e4e158 100644 --- a/src/plugins/intel_gpu/src/graph/impls/cpu/non_max_suppression.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/cpu/non_max_suppression.cpp @@ -6,7 +6,7 @@ #include "primitive_inst.h" #include "register.hpp" #include "cpu_impl_helpers.hpp" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include #include diff --git a/src/plugins/intel_gpu/src/graph/impls/cpu/proposal.cpp b/src/plugins/intel_gpu/src/graph/impls/cpu/proposal.cpp index bd74efe4ee0..6fe3fac8676 100644 --- a/src/plugins/intel_gpu/src/graph/impls/cpu/proposal.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/cpu/proposal.cpp @@ -4,7 +4,7 @@ #include "proposal_inst.h" #include "intel_gpu/runtime/engine.hpp" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "intel_gpu/runtime/error_handler.hpp" #include "register.hpp" diff --git a/src/plugins/intel_gpu/src/graph/impls/cpu/read_value.cpp b/src/plugins/intel_gpu/src/graph/impls/cpu/read_value.cpp index bbecd0eb48d..2b85f61d3ad 100644 --- a/src/plugins/intel_gpu/src/graph/impls/cpu/read_value.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/cpu/read_value.cpp @@ -3,8 +3,9 @@ // #include "read_value_inst.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "register.hpp" +#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace cpu { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/activation.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/activation.cpp index bda9bfd5696..4f219b18a34 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/activation.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/activation.cpp @@ -2,16 +2,14 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "activation/activation_kernel_base.h" -#include "activation/activation_kernel_selector.h" -#include "activation_inst.h" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" #include "primitive_base.hpp" +#include "activation_inst.h" +#include "activation/activation_kernel_base.h" +#include "activation/activation_kernel_selector.h" + namespace { -inline void convert_new_activation_func(const activation& prim, std::vector& params) { +inline void convert_new_activation_func(const cldnn::activation& prim, std::vector& params) { params.insert(params.begin(), {get_kernel_selector_activation_param(prim.activation_function), prim.additional_params.a, prim.additional_params.b}); diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/adaptive_pooling.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/adaptive_pooling.cpp index f069d4aa7c1..768cb82fb2c 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/adaptive_pooling.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/adaptive_pooling.cpp @@ -2,14 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "intel_gpu/primitives/adaptive_pooling.hpp" +#include "primitive_base.hpp" +#include "adaptive_pooling_inst.h" #include "adaptive_pooling/adaptive_pooling_kernel_ref.h" #include "adaptive_pooling/adaptive_pooling_kernel_selector.h" -#include "adaptive_pooling_inst.h" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" -#include "primitive_base.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/arg_max_min.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/arg_max_min.cpp index 3d78a09e8fa..98a6cf168d4 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/arg_max_min.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/arg_max_min.cpp @@ -2,11 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "arg_max_min_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "arg_max_min_inst.h" #include "arg_max_min/arg_max_min_kernel_selector.h" #include "arg_max_min/arg_max_min_kernel_base.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/batch_to_space.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/batch_to_space.cpp index ec7e12d2242..4d4aa93dfdd 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/batch_to_space.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/batch_to_space.cpp @@ -2,17 +2,10 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "batch_to_space_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" +#include "batch_to_space_inst.h" #include "batch_to_space/batch_to_space_kernel_selector.h" #include "batch_to_space/batch_to_space_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" -#include "data_inst.h" -#include - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/binary_convolution.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/binary_convolution.cpp index 4a80568a948..4d91b87fc03 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/binary_convolution.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/binary_convolution.cpp @@ -2,16 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "intel_gpu/primitives/quantize.hpp" -#include "binary_convolution_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "binary_convolution_inst.h" #include "kernel_selector/kernels/binary_convolution/binary_convolution_kernel_selector.h" #include "kernel_selector/kernels/binary_convolution/binary_convolution_params.h" -#include -#include namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/border.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/border.cpp index 282c54a1439..f3f83f3a016 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/border.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/border.cpp @@ -2,14 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // +#include "primitive_base.hpp" + #include "border_inst.h" -#include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" #include "border/border_kernel_selector.h" #include "border/border_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/broadcast.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/broadcast.cpp index 662dbf1ceb1..e2f785a8d6f 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/broadcast.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/broadcast.cpp @@ -2,14 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "broadcast_inst.h" - #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "broadcast_inst.h" #include "broadcast/broadcast_kernel_selector.h" #include "broadcast/broadcast_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/bucketize.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/bucketize.cpp index f6ee0585c08..2cc508e8e0f 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/bucketize.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/bucketize.cpp @@ -2,11 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // +#include "primitive_base.hpp" + +#include "bucketize_inst.hpp" #include "bucketize/bucketize_kernel_ref.hpp" #include "bucketize/bucketize_kernel_selector.hpp" -#include "bucketize_inst.hpp" -#include "impls/implementation_map.hpp" -#include "primitive_base.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/concatenation.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/concatenation.cpp index f766353d9a2..16e966830b1 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/concatenation.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/concatenation.cpp @@ -2,16 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "concatenation_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "concatenation_inst.h" #include "concatenation/concatenation_kernel_selector.h" #include "concatenation/concatenation_kernel_base.h" -#include - namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/convert_color.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/convert_color.cpp index 7fc4b91cb66..66c9088c707 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/convert_color.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/convert_color.cpp @@ -2,17 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "convert_color_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "convert_color_inst.h" #include "convert_color/convert_color_kernel_selector.h" #include "convert_color/convert_color_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" -#include "data_inst.h" -#include - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/convolution.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/convolution.cpp index ede4bffa479..ec822fd2a35 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/convolution.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/convolution.cpp @@ -2,16 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "convolution_inst.h" -#include "eltwise_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "convolution_inst.h" #include "convolution/convolution_kernel_selector.h" #include "convolution/convolution_params.h" -#include -#include namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/crop.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/crop.cpp index acf400f3b7e..760e7d6ba4a 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/crop.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/crop.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "crop_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "crop_inst.h" #include "eltwise/eltwise_kernel_selector.h" #include "eltwise/eltwise_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/ctc_greedy_decoder.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/ctc_greedy_decoder.cpp index 7cc15cefaf8..8f931a779eb 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/ctc_greedy_decoder.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/ctc_greedy_decoder.cpp @@ -2,18 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "ctc_greedy_decoder_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "ctc_greedy_decoder_inst.h" #include "ctc_greedy_decoder/ctc_greedy_decoder_kernel_selector.h" #include "ctc_greedy_decoder/ctc_greedy_decoder_kernel_base.h" -#include - -using namespace cldnn; - namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/ctc_loss.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/ctc_loss.cpp index 2635c327b72..78c6caf9ca5 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/ctc_loss.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/ctc_loss.cpp @@ -2,11 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // +#include "primitive_base.hpp" + +#include "ctc_loss_inst.hpp" #include "ctc_loss/ctc_loss_kernel_ref.hpp" #include "ctc_loss/ctc_loss_kernel_selector.hpp" -#include "ctc_loss_inst.hpp" -#include "impls/implementation_map.hpp" -#include "primitive_base.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/cum_sum.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/cum_sum.cpp index 3e2e102db6d..f3b95dc218d 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/cum_sum.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/cum_sum.cpp @@ -2,15 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "cum_sum_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "cum_sum_inst.h" #include "cum_sum/cum_sum_kernel_selector.h" #include "cum_sum/cum_sum_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/custom_primitive.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/custom_primitive.cpp index 5c0e7e6f49a..c4a901b72e8 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/custom_primitive.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/custom_primitive.cpp @@ -2,13 +2,10 @@ // SPDX-License-Identifier: Apache-2.0 // +#include "primitive_base.hpp" + #include "custom_gpu_primitive_inst.h" -#include "intel_gpu/runtime/engine.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" #include "jitter.h" -#include "intel_gpu/runtime/error_handler.hpp" -#include "register.hpp" #include #include @@ -142,9 +139,7 @@ static void add_layout_to_jit(kernel_selector::jit_constants& mem_consts, const {data_types::f32, "float"}, }; - if (dataTypeToIndex.find(l.data_type) == dataTypeToIndex.end()) { - CLDNN_ERROR_MESSAGE("add layout to jit", "Unhandled data type in layout"); - } + OPENVINO_ASSERT(dataTypeToIndex.find(l.data_type) != dataTypeToIndex.end(), "[GPU] Add layout to jit error: unhandled data type in layout"); mem_consts.AddConstant(kernel_selector::MakeJitConstant(name + "_TYPE", dataTypeToIndex.at(l.data_type))); diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/deconvolution.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/deconvolution.cpp index 9297ebb5f6b..d25826c63ca 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/deconvolution.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/deconvolution.cpp @@ -2,14 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "deconvolution_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "deconvolution_inst.h" #include "deconvolution/deconvolution_kernel_selector.h" #include "deconvolution/deconvolution_kernel_base.h" -#include namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/deformable_convolution.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/deformable_convolution.cpp index 0531dd2f0fb..048a83f3cec 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/deformable_convolution.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/deformable_convolution.cpp @@ -2,14 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "deformable_convolution_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "deformable_convolution_inst.h" #include "convolution/convolution_kernel_selector.h" #include "convolution/convolution_params.h" -#include namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/depth_to_space.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/depth_to_space.cpp index d65273f216a..14e1eb0711f 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/depth_to_space.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/depth_to_space.cpp @@ -2,16 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "depth_to_space_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "depth_to_space_inst.h" #include "depth_to_space/depth_to_space_kernel_selector.h" #include "depth_to_space/depth_to_space_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" -#include "common_types.h" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/detection_output.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/detection_output.cpp index 205193ea3f5..84978c2d5c1 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/detection_output.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/detection_output.cpp @@ -2,14 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "detection_output_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "detection_output_inst.h" #include "detection_output/detection_output_kernel_selector.h" #include "detection_output/detection_output_kernel_ref.h" -#include namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/dft.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/dft.cpp index 387e78d092c..e5cd0f77a5b 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/dft.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/dft.cpp @@ -2,16 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include -#include -#include -#include - -#include -#include - #include "primitive_base.hpp" +#include "dft_inst.h" +#include "dft/dft_kernel_ref.h" +#include "dft/dft_kernel_selector.h" + namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/eltwise.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/eltwise.cpp index 27ae5bf50d9..64418f5c2e6 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/eltwise.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/eltwise.cpp @@ -2,14 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "eltwise_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "eltwise_inst.h" #include "eltwise/eltwise_kernel_selector.h" #include "eltwise/eltwise_kernel_base.h" -#include namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/embedding_bag.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/embedding_bag.cpp index eec9f0a1e20..0d4703c3359 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/embedding_bag.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/embedding_bag.cpp @@ -2,16 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "embedding_bag_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "embedding_bag_inst.h" #include "embedding_bag/embedding_bag_kernel_selector.h" #include "embedding_bag/embedding_bag_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" -#include "data_inst.h" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_detection_output.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_detection_output.cpp index d905fcc3026..0806d38e666 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_detection_output.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_detection_output.cpp @@ -2,12 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // +#include "primitive_base.hpp" + +#include "experimental_detectron_detection_output_inst.hpp" #include "ed_do/detection_output_kernel_ref.h" #include "ed_do/detection_output_kernel_selector.h" -#include "experimental_detectron_detection_output_inst.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" -#include "primitive_base.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_generate_proposals_single_image.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_generate_proposals_single_image.cpp index e0832fcb03b..29d4b0aea1c 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_generate_proposals_single_image.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_generate_proposals_single_image.cpp @@ -2,14 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "experimental_detectron_generate_proposals_single_image_inst.hpp" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "experimental_detectron_generate_proposals_single_image_inst.hpp" #include "ed_gpsi/generate_proposals_single_image_kernel_selector.h" #include "ed_gpsi/generate_proposals_single_image_kernel_ref.h" - namespace cldnn { namespace ocl { struct experimental_detectron_generate_proposals_single_image_impl 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 index 1c989a3f234..d5a720a5088 100644 --- 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 @@ -2,18 +2,12 @@ // 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" +#include "experimental_detectron_prior_grid_generator_inst.h" +#include "ed_pgg/prior_grid_generator_kernel_ref.h" +#include "ed_pgg/prior_grid_generator_kernel_selector.h" + namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_roi_feature_extractor.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_roi_feature_extractor.cpp index 257e0607254..0c15c41fe33 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_roi_feature_extractor.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_roi_feature_extractor.cpp @@ -2,11 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "intel_gpu/primitives/experimental_detectron_roi_feature_extractor.hpp" -#include "experimental_detectron_roi_feature_extractor_inst.hpp" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "experimental_detectron_roi_feature_extractor_inst.hpp" #include "ed_rfe/roi_feature_extractor_kernel_selector.h" #include "ed_rfe/roi_feature_extractor_kernel_ref.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_topk_rois.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_topk_rois.cpp index da6613a5631..0211153bc68 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_topk_rois.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/experimental_detectron_topk_rois.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include -#include "intel_gpu/runtime/error_handler.hpp" -#include -#include -#include #include "primitive_base.hpp" -#include + +#include "experimental_detectron_topk_rois_inst.h" +#include "ed_topkroi/topk_rois_ref.h" +#include "ed_topkroi/topk_rois_kernel_selector.h" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/extract_image_patches.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/extract_image_patches.cpp index 830644dbd60..839b4bfda60 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/extract_image_patches.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/extract_image_patches.cpp @@ -2,12 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "extract_image_patches_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" +#include "extract_image_patches_inst.h" #include "extract_image_patches/extract_image_patches_kernel_selector.h" #include "extract_image_patches/extract_image_patches_kernel_ref.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/eye.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/eye.cpp index 6c3470a8e51..c6d2d279938 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/eye.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/eye.cpp @@ -2,19 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include -#include -#include -#include - -#include -#include -#include -#include -#include - #include "primitive_base.hpp" +#include "eye_inst.h" +#include "eye/eye_kernel_ref.h" +#include "eye/eye_kernel_selector.h" + namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/fully_connected.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/fully_connected.cpp index 12ba395626c..c6e703d63ff 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/fully_connected.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/fully_connected.cpp @@ -2,20 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "fully_connected_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "fully_connected_inst.h" #include "fully_connected/fully_connected_kernel_selector.h" #include "fully_connected/fully_connected_params.h" -#include "intel_gpu/runtime/error_handler.hpp" - -#include "intel_gpu/primitives/reorder.hpp" -#include "intel_gpu/primitives/input_layout.hpp" -#include -#include - namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/gather.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/gather.cpp index 2c5037ec053..27348e0b06d 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/gather.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/gather.cpp @@ -2,15 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "gather_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "gather_inst.h" #include "gather/gather_kernel_selector.h" #include "gather/gather_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/gather_elements.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/gather_elements.cpp index 9697f666472..6eb050e9113 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/gather_elements.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/gather_elements.cpp @@ -2,15 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "gather_elements_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "gather_elements_inst.h" #include "gather/gather_elements_kernel_selector.h" #include "gather/gather_elements_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/gather_nd.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/gather_nd.cpp index 67ca4845321..3346e7e4317 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/gather_nd.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/gather_nd.cpp @@ -2,15 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "gather_nd_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "gather_nd_inst.h" #include "gather/gather_nd_kernel_selector.h" #include "gather/gather_nd_kernel_ref.h" -using namespace cldnn; - namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/gather_tree.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/gather_tree.cpp index 907c8ecb38e..3c80326b2ce 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/gather_tree.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/gather_tree.cpp @@ -2,14 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "gather_tree_inst.h" - #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "gather_tree_inst.h" #include "gather_tree/gather_tree_kernel_selector.h" #include "gather_tree/gather_tree_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/gemm.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/gemm.cpp index d475fd328dc..911445a3dfa 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/gemm.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/gemm.cpp @@ -2,14 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // +#include "primitive_base.hpp" + +#include "gemm_inst.h" #include "gemm/gemm_kernel_base.h" #include "gemm/gemm_kernel_selector.h" -#include "gemm_inst.h" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include -#include "kernel_selector_helper.h" -#include "primitive_base.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/generate_proposals.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/generate_proposals.cpp index 29359c20bb3..46928f6d983 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/generate_proposals.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/generate_proposals.cpp @@ -2,15 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "generate_proposals_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "generate_proposals_inst.h" #include "generate_proposals/generate_proposals_kernel_selector.h" #include "generate_proposals/generate_proposals_kernel_ref.h" - - namespace cldnn { namespace ocl { struct generate_proposals_impl diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/generic_layer.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/generic_layer.cpp index ed7b09fe6bd..24f5464a7f7 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/generic_layer.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/generic_layer.cpp @@ -2,12 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // +#include "primitive_base.hpp" + #include "generic_layer_inst.h" -#include "intel_gpu/runtime/engine.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" -#include "register.hpp" -#include namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/grid_sample.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/grid_sample.cpp index 0063a1c6003..5934f7d30c6 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/grid_sample.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/grid_sample.cpp @@ -2,11 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // +#include "primitive_base.hpp" + +#include "grid_sample_inst.hpp" #include "grid_sample/grid_sample_kernel_ref.hpp" #include "grid_sample/grid_sample_kernel_selector.hpp" -#include "grid_sample_inst.hpp" -#include "impls/implementation_map.hpp" -#include "primitive_base.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/grn.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/grn.cpp index 8a6b5beea74..543502aa933 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/grn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/grn.cpp @@ -2,18 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "grn_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "grn_inst.h" #include "grn/grn_kernel_selector.h" #include "grn/grn_kernel_base.h" -#include - -using namespace cldnn; - namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/kernel_selector_helper.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/kernel_selector_helper.cpp similarity index 99% rename from src/plugins/intel_gpu/src/graph/kernel_selector_helper.cpp rename to src/plugins/intel_gpu/src/graph/impls/ocl/kernel_selector_helper.cpp index 1ebbed40baa..8bfeff2a57c 100644 --- a/src/plugins/intel_gpu/src/graph/kernel_selector_helper.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/kernel_selector_helper.cpp @@ -34,7 +34,10 @@ #include #include + namespace { +using namespace cldnn; + kernel_selector::dev_type get_device_type(cldnn::device_type type) { switch (type) { case cldnn::device_type::integrated_gpu: @@ -47,6 +50,8 @@ kernel_selector::dev_type get_device_type(cldnn::device_type type) { } } // namespace +namespace cldnn { + kernel_selector::data_type to_data_type(data_types dt) { switch (dt) { case cldnn::data_types::bin: @@ -1244,3 +1249,5 @@ void kernel_impl_params::load(BinaryInputBuffer& ib) { #endif // ENABLE_ONEDNN_FOR_GPU ib >> primary_input_idx; } + +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/graph/include/kernel_selector_helper.h b/src/plugins/intel_gpu/src/graph/impls/ocl/kernel_selector_helper.h similarity index 72% rename from src/plugins/intel_gpu/src/graph/include/kernel_selector_helper.h rename to src/plugins/intel_gpu/src/graph/impls/ocl/kernel_selector_helper.h index 4112073b1af..7057006c147 100644 --- a/src/plugins/intel_gpu/src/graph/include/kernel_selector_helper.h +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/kernel_selector_helper.h @@ -8,7 +8,6 @@ #include "intel_gpu/runtime/engine.hpp" #include "intel_gpu/runtime/utils.hpp" #include "intel_gpu/runtime/tensor.hpp" -#include "intel_gpu/runtime/error_handler.hpp" #include "intel_gpu/primitives/eltwise.hpp" #include "intel_gpu/primitives/quantize.hpp" #include "intel_gpu/primitives/activation.hpp" @@ -16,6 +15,7 @@ #include "kernel_selector_params.h" #include "kernel_selector_common.h" +#include "kernel_impl_params.hpp" #include "tensor_type.h" #include "fused_primitive_desc.h" @@ -24,16 +24,6 @@ #include #include -using namespace cldnn; - -namespace cldnn { -enum class data_types : size_t; -struct format; -struct layout; -struct program; -struct fused_primitive_desc; -} // namespace cldnn - namespace kernel_selector { using n_dims = kernel_selector::Tensor::NDims; using kernel_data = kernel_selector::KernelData; @@ -91,6 +81,13 @@ using generic_kernel_params = kernel_selector::GenericKernelParams; } // namespace kernel_selector +namespace cldnn { +enum class data_types : size_t; +struct format; +struct layout; +struct program; +struct fused_primitive_desc; + kernel_selector::data_type to_data_type(data_types dt); data_types from_data_type(kernel_selector::data_type dt); kernel_selector::weights_type to_weights_type(data_types dt); @@ -104,100 +101,6 @@ kernel_selector::weights_tensor convert_weights_tensor(const layout& l, bool is_ layout from_weights_tensor(const kernel_selector::weights_tensor& t); kernel_selector::activation_function get_kernel_selector_activation_param(activation_func activation_func); -struct kernel_impl_params { - bool has_runtime_layouts = false; - const program *prog; - std::shared_ptr desc; - size_t unique_id; - std::vector input_layouts; - std::vector output_layouts; - std::vector input_offsets; - std::vector fused_desc; -#ifdef ENABLE_ONEDNN_FOR_GPU - std::vector fused_desc_onednn; -#endif // ENABLE_ONEDNN_FOR_GPU - - optional_layout weights_layout = optional_layout(); - - optional_layout bias_layout = optional_layout(); - optional_layout weights_zero_points_layout = optional_layout(); - optional_layout activations_zero_points_layout = optional_layout(); - optional_layout compensation_layout = optional_layout(); - - std::map memory_deps = {}; - size_t primary_input_idx = 0; - - memory::ptr reordered_weights = nullptr; - - kernel_impl_params() {} - - kernel_impl_params(program& _prog, - std::shared_ptr _desc, - size_t _uid, - const std::vector& _in_layouts, - const std::vector& _out_layouts, - const std::vector& _fused_descs) - : has_runtime_layouts(true) - , prog(&_prog) - , desc(_desc) - , unique_id(_uid) - , input_layouts(_in_layouts) - , output_layouts(_out_layouts) - , fused_desc(_fused_descs) - , primary_input_idx(0) { - } - - layout get_input_layout(size_t idx = 0) const { - OPENVINO_ASSERT(input_layouts.size() > idx, - "The size of input layouts must be greater than the requested index: ", - "Requested index is ", idx, ", ", - "but the size of input layouts is ", input_layouts.size()); - return input_layouts[idx]; - } - - layout get_non_padded_input_layout(size_t idx = 0) const { - auto input_layout = get_input_layout(idx); - auto result = layout({input_layout.get_partial_shape(), input_layout.data_type, input_layout.format}); - return result; - } - - layout get_output_layout(size_t idx = 0) const { - OPENVINO_ASSERT(output_layouts.size() > idx, - "The size of output layouts must be greater than the requested index: ", - "Requested index is ", idx, ",", - "but the size of output layouts is ", output_layouts.size()); - return output_layouts[idx]; - } - - bool has_fused_primitives() const { return !fused_desc.empty(); } - - layout get_fused_output_layout() const { - if (fused_desc.empty()) - return layout(data_types::f32, format::bfyx, tensor()); - return fused_desc.back().output_layout; - } - - bool is_dynamic() const { - for (auto i : input_layouts) - if (i.is_dynamic()) - return true; - for (auto i : output_layouts) - if (i.is_dynamic()) - return true; - return false; - } - - template - std::shared_ptr typed_desc() const { return std::static_pointer_cast(desc); } - - void save(BinaryOutputBuffer& ob) const; - void load(BinaryInputBuffer& ib); - const program& get_program() const { - OPENVINO_ASSERT(prog != nullptr, "[GPU] Program pointer in kernel_impl_params in not initialized"); - return *prog; - } -}; - template kernel_selector::dim_tensor convert_dim_vector(const tensor& t) { const auto& sizes = t.sizes(format::bfwzyx); @@ -239,10 +142,8 @@ inline params_t get_default_params(const kernel_impl_params& param_info, bool is kernel_selector::fused_operation_desc desc; desc.op_params = std::move(fused_prim.f_param); - if (!desc.op_params) { - CLDNN_ERROR_MESSAGE(param_info.desc->id, "Invalid fused operation (" + param_info.desc->id + ") of type " + - param_info.desc->type_string()); - } + OPENVINO_ASSERT(desc.op_params != nullptr, "[GPU] Invalid fused operation (", param_info.desc->id , ") of type ", param_info.desc->type_string()); + desc.dep_idx_start = fused_prim.dep_start_idx; desc.dep_size = fused_prim.deps.size(); @@ -337,3 +238,4 @@ template inline optional_params_t get_default_weights_bias_optional_params(const program& program) { return get_default_optional_params(program); } +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/lrn.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/lrn.cpp index 4b770ab83a2..0e556c7394c 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/lrn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/lrn.cpp @@ -2,11 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "lrn_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "lrn_inst.h" #include "lrn/lrn_kernel_selector.h" #include "lrn/lrn_kernel_base.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_dynamic_input.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_dynamic_input.cpp index b30d68944b7..fc21a1bcaa8 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_dynamic_input.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_dynamic_input.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "lstm_dynamic_input_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "lstm_dynamic_input_inst.h" #include "lstm_dynamic/lstm_dynamic_input_kernel_selector.h" #include "lstm_dynamic/lstm_dynamic_input_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_dynamic_timeloop.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_dynamic_timeloop.cpp index 1e3935e65fa..a7812a02973 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_dynamic_timeloop.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_dynamic_timeloop.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "lstm_dynamic_timeloop_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "lstm_dynamic_timeloop_inst.h" #include "lstm_dynamic/lstm_dynamic_timeloop_kernel_selector.h" #include "lstm_dynamic/lstm_dynamic_timeloop_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_elt.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_elt.cpp index 33353d3dcc3..bfaf79cec78 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_elt.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_elt.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "lstm_elt_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "lstm_elt_inst.h" #include "lstm/lstm_elt_kernel_selector.h" #include "lstm/lstm_elt_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { @@ -36,60 +34,49 @@ protected: } public: - static std::unique_ptr create(const lstm_elt_node& arg, const kernel_impl_params& impl_param) { - const auto& prim = arg.get_primitive(); - auto lstm_elt_params = get_default_params(impl_param); - auto lstm_elt_optional_params = - get_default_optional_params(arg.get_program()); + static kernel_params_t get_kernel_params(const kernel_impl_params& impl_param) { + const auto& primitive = impl_param.typed_desc(); + auto params = get_default_params(impl_param); + auto optional_params = get_default_optional_params(impl_param.get_program()); - if (arg.cell_term()) { + if (!primitive->cell.empty()) { const auto& cell_idx = 1; const auto& cell_layout = impl_param.input_layouts[cell_idx]; - lstm_elt_params.SetCell(convert_data_tensor(cell_layout)); + params.SetCell(convert_data_tensor(cell_layout)); // TODO: make a generic function to get the direction if (cell_layout.spatial(1) > 1) { - lstm_elt_params.cell_direction = arg.direction(); + params.cell_direction = primitive->direction; } } - if (!prim->activations.empty()) { - auto a_sz = prim->activations.size(); - auto param_sz = prim->activation_params.size(); - if (param_sz) { - CLDNN_ERROR_NOT_EQUAL(arg.id(), - "number of activations", - a_sz, - "number of activation parameters", - param_sz, - "activations/parameters num mismatch"); - } + if (!primitive->activations.empty()) { + auto a_sz = primitive->activations.size(); + auto param_sz = primitive->activation_params.size(); + OPENVINO_ASSERT(param_sz == 0|| a_sz == param_sz, "[GPU] Unexpected activation params count in lstm_elt impl: ", param_sz); for (size_t i = 0; i < a_sz; i++) { - lstm_elt_params.activations.emplace_back(get_kernel_selector_activation_param(prim->activations[i]), - param_sz ? prim->activation_params[i].a : 0.0f, - param_sz ? prim->activation_params[i].b : 0.0f); + params.activations.emplace_back(get_kernel_selector_activation_param(primitive->activations[i]), + param_sz ? primitive->activation_params[i].a : 0.0f, + param_sz ? primitive->activation_params[i].b : 0.0f); } } - if (prim->clip > 0.0f) { - lstm_elt_params.activations.emplace_back(get_kernel_selector_activation_param(activation_func::clamp), -prim->clip, prim->clip); + if (primitive->clip > 0.0f) { + params.activations.emplace_back(get_kernel_selector_activation_param(activation_func::clamp), -primitive->clip, primitive->clip); } - lstm_elt_params.SetOffsetOrder(static_cast(arg.offset_order())); - lstm_elt_params.clip = arg.clip(); - lstm_elt_params.input_forget = arg.input_forget(); - lstm_elt_params.direction = arg.direction(); + params.SetOffsetOrder(static_cast(primitive->offset_order)); + params.clip = primitive->clip; + params.input_forget = primitive->input_forget; + params.direction = primitive->direction; - auto& kernel_selector = kernel_selector::lstm_elt_kernel_selector::Instance(); - auto best_kernel = kernel_selector.get_best_kernel(lstm_elt_params, lstm_elt_optional_params); - - return make_unique(best_kernel); + return {params, optional_params}; } }; namespace detail { attach_lstm_elt_impl::attach_lstm_elt_impl() { - implementation_map::add(impl_types::ocl, lstm_elt_impl::create, { + implementation_map::add(impl_types::ocl, typed_primitive_impl_ocl::create, { std::make_tuple(data_types::f32, format::bfyx), std::make_tuple(data_types::f16, format::bfyx), std::make_tuple(data_types::f32, format::fyxb), diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_gemm.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_gemm.cpp index 6c6c3eaa8f9..94f20b7d42a 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_gemm.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/lstm_gemm.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "lstm_gemm_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "lstm_gemm_inst.h" #include "lstm/lstm_gemm_kernel_selector.h" #include "lstm/lstm_gemm_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/matrix_nms.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/matrix_nms.cpp index dd7f0b3d9ca..e3744e21a3e 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/matrix_nms.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/matrix_nms.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "data_inst.h" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" +#include "primitive_base.hpp" + +#include "matrix_nms_inst.h" #include "matrix_nms/matrix_nms_kernel_ref.h" #include "matrix_nms/matrix_nms_kernel_selector.h" -#include "matrix_nms_inst.h" -#include "primitive_base.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/multiclass_nms.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/multiclass_nms.cpp index 7752d5cf538..f1976a7e0ad 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/multiclass_nms.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/multiclass_nms.cpp @@ -2,14 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "multiclass_nms_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "multiclass_nms_inst.h" #include "multiclass_nms/multiclass_nms_kernel_ref.h" #include "multiclass_nms/multiclass_nms_kernel_selector.h" - namespace cldnn { namespace ocl { namespace { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/mutable_data.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/mutable_data.cpp index f6da456fb5c..9d6f76f1151 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/mutable_data.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/mutable_data.cpp @@ -2,9 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "mutable_data_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" + +#include "mutable_data_inst.h" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/mvn.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/mvn.cpp index a6ca592a7d9..1223f732157 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/mvn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/mvn.cpp @@ -2,18 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "mvn_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "mvn_inst.h" #include "mvn/mvn_kernel_selector.h" #include "mvn/mvn_kernel_base.h" -#include - -using namespace cldnn; - namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/non_max_suppression.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/non_max_suppression.cpp index 6955fd86a78..a8cff3afe87 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/non_max_suppression.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/non_max_suppression.cpp @@ -2,13 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // +#include "primitive_base.hpp" + +#include "non_max_suppression_inst.h" #include "data_inst.h" -#include "kernel_selector_helper.h" #include "non_max_suppression/non_max_suppression_kernel_ref.h" #include "non_max_suppression/non_max_suppression_kernel_selector.h" -#include "non_max_suppression_inst.h" -#include "primitive_base.hpp" -#include "impls/implementation_map.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/non_zero.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/non_zero.cpp index 9103d1aec2e..e796cec93e0 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/non_zero.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/non_zero.cpp @@ -2,17 +2,13 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "non_zero_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "non_zero_inst.h" #include "non_zero/count_nonzero_kernel_ref.h" #include "non_zero/count_nonzero_kernel_selector.h" #include "non_zero/gather_nonzero_kernel_ref.h" #include "non_zero/gather_nonzero_kernel_selector.h" -#include "intel_gpu/runtime/error_handler.hpp" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/normalize.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/normalize.cpp index e1fa378b435..26eaff63962 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/normalize.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/normalize.cpp @@ -2,18 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "normalize_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "normalize_inst.h" #include "normalize/normalize_kernel_selector.h" #include "normalize/normalize_kernel_base.h" -#include - -using namespace cldnn; - namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/one_hot.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/one_hot.cpp index ea6c881121f..b2d6021575b 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/one_hot.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/one_hot.cpp @@ -2,15 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "one_hot_inst.h" - #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "one_hot_inst.h" #include "one_hot/one_hot_kernel_selector.h" #include "one_hot/one_hot_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" -#include namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/permute.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/permute.cpp index d353f879dcb..3f8842d1e5b 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/permute.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/permute.cpp @@ -2,16 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "permute_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "permute_inst.h" #include "permute/permute_kernel_selector.h" #include "permute/permute_kernel_ref.h" -using namespace cldnn; - namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/pooling.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/pooling.cpp index 28b06faf37f..de63adf0c27 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/pooling.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/pooling.cpp @@ -2,15 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "pooling_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "pooling_inst.h" #include "pooling/pooling_kernel_selector.h" #include "pooling/pooling_kernel_base.h" #include "ngraph/validation_util.hpp" -#include namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/primitive_base.hpp b/src/plugins/intel_gpu/src/graph/impls/ocl/primitive_base.hpp index 10c1ee1cf40..e247c981f03 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/primitive_base.hpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/primitive_base.hpp @@ -4,11 +4,6 @@ #pragma once -#include -#include "primitive_inst.h" -#include "intel_gpu/graph/program.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" #include "intel_gpu/graph/network.hpp" #include "intel_gpu/graph/serialization/binary_buffer.hpp" #include "intel_gpu/graph/serialization/cl_kernel_data_serializer.hpp" @@ -16,7 +11,13 @@ #include "intel_gpu/graph/serialization/set_serializer.hpp" #include "intel_gpu/graph/serialization/string_serializer.hpp" #include "intel_gpu/graph/serialization/vector_serializer.hpp" +#include "intel_gpu/graph/program.hpp" + +#include "primitive_inst.h" +#include "kernel_selector_helper.h" #include "register.hpp" +#include "implementation_map.hpp" + #include #include #include diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/prior_box.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/prior_box.cpp index 7c6e911309a..aeaf96f98de 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/prior_box.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/prior_box.cpp @@ -2,16 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include -#include -#include - -#include -#include - -#include "intel_gpu/runtime/error_handler.hpp" #include "primitive_base.hpp" +#include "prior_box_inst.h" +#include "prior_box/prior_box_kernel_ref.h" +#include "prior_box/prior_box_kernel_selector.h" + namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/pyramid_roi_align.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/pyramid_roi_align.cpp index 039ef0adcde..2f95f8268d7 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/pyramid_roi_align.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/pyramid_roi_align.cpp @@ -3,12 +3,10 @@ // #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "pyramid_roi_align_inst.h" #include "pyramid_roi_align/pyramid_roi_align_kernel_selector.h" #include "pyramid_roi_align/pyramid_roi_align_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" -#include "pyramid_roi_align_inst.h" #include diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/quantize.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/quantize.cpp index a500f0357e0..c185c7ac978 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/quantize.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/quantize.cpp @@ -2,15 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "quantize_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "quantize_inst.h" #include "quantize/quantize_kernel_selector.h" #include "quantize/quantize_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/random_uniform.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/random_uniform.cpp index 7fd2326edd3..ce81cddb144 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/random_uniform.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/random_uniform.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include -#include -#include "intel_gpu/runtime/error_handler.hpp" -#include -#include #include "primitive_base.hpp" -#include + +#include "random_uniform_inst.h" +#include "random_uniform/random_uniform_kernel_ref.h" +#include "random_uniform/random_uniform_kernel_selector.h" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/range.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/range.cpp index 6e8c4380366..959990e62d1 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/range.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/range.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include #include "primitive_base.hpp" -#include -#include -#include -#include -#include + +#include "range_inst.h" +#include "range/range_kernel_selector.h" +#include "range/range_kernel_ref.h" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/reduce.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/reduce.cpp index f64ac4c5095..f7d8ba3f96c 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/reduce.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/reduce.cpp @@ -2,17 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "reduce_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "reduce_inst.h" #include "reduce/reduce_kernel_selector.h" #include "reduce/reduce_kernel_ref.h" -#include "reduce/reduce_kernel_b_fs_yx_fsv16.h" -#include "intel_gpu/runtime/error_handler.hpp" -#include "data_inst.h" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/region_yolo.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/region_yolo.cpp index b0078b97116..698816747f7 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/region_yolo.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/region_yolo.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "region_yolo_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "region_yolo_inst.h" #include "region_yolo/region_yolo_kernel_selector.h" #include "region_yolo/region_yolo_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/reorder.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/reorder.cpp index 96748ba4a94..7528c7c178c 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/reorder.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/reorder.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "reorder_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "reorder_inst.h" #include "reorder/reorder_kernel_selector.h" #include "reorder/reorder_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/reorg_yolo.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/reorg_yolo.cpp index 257685fbf98..275da00223f 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/reorg_yolo.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/reorg_yolo.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "reorg_yolo_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "reorg_yolo_inst.h" #include "reorg_yolo/reorg_yolo_kernel_selector.h" #include "reorg_yolo/reorg_yolo_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/resample.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/resample.cpp index 0904ed8288b..c01d6bc6f2b 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/resample.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/resample.cpp @@ -2,16 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include +#include "primitive_base.hpp" #include "resample_inst.h" -#include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" #include "kernel_selector/kernels/resample/resample_kernel_selector.h" #include "kernel_selector/kernels/resample/resample_kernel_base.h" -#include "intel_gpu/runtime/half.hpp" +#include namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/reshape.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/reshape.cpp index def35a6e78a..19a3c90e8dc 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/reshape.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/reshape.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "reshape_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "reshape_inst.h" #include "reshape/reshape_kernel_ref.h" #include "reshape/reshape_kernel_selector.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/reverse.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/reverse.cpp index 2d6b2601ac6..f1d8132a68a 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/reverse.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/reverse.cpp @@ -1,16 +1,12 @@ -// Copyright (C) 2022 Intel Corporation +// Copyright (C) 2022-2023 Intel Corporation // SPDX-License-Identifier: Apache-2.0 // -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" #include "primitive_base.hpp" + +#include "reverse_inst.h" #include "reverse/reverse_kernel_ref.h" #include "reverse/reverse_kernel_selector.h" -#include "reverse_inst.h" -#include "impls/implementation_map.hpp" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/reverse_sequence.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/reverse_sequence.cpp index a51973e0bd3..f44985e4c08 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/reverse_sequence.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/reverse_sequence.cpp @@ -2,15 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "reverse_sequence_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "reverse_sequence_inst.h" #include "reverse_sequence/reverse_sequence_kernel_selector.h" #include "reverse_sequence/reverse_sequence_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/roi_align.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/roi_align.cpp index c17b2e9d97a..c8ca8d86a35 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/roi_align.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/roi_align.cpp @@ -3,12 +3,10 @@ // #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "roi_align_inst.h" #include "roi_align/roi_align_kernel_ref.h" #include "roi_align/roi_align_kernel_selector.h" -#include "roi_align_inst.h" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/roi_pooling.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/roi_pooling.cpp index d99c92e4a95..d3d7ff72489 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/roi_pooling.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/roi_pooling.cpp @@ -2,11 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "roi_pooling_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "roi_pooling_inst.h" #include "roi_pooling/roi_pooling_kernel_selector.h" #include "roi_pooling/roi_pooling_kernel_ref.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/roll.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/roll.cpp index 3d2574d836c..5981cfc963d 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/roll.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/roll.cpp @@ -1,11 +1,12 @@ -// Copyright (C) 2022 Intel Corporation +// Copyright (C) 2022-2023 Intel Corporation // SPDX-License-Identifier: Apache-2.0 // -#include "impls/implementation_map.hpp" + #include "primitive_base.hpp" + +#include "roll_inst.hpp" #include "roll/roll_kernel_ref.hpp" #include "roll/roll_kernel_selector.hpp" -#include "roll_inst.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_elements_update.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_elements_update.cpp index e95353de1bb..54f54717b1b 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_elements_update.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_elements_update.cpp @@ -2,15 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "scatter_elements_update_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "scatter_elements_update_inst.h" #include "scatter_update/scatter_elements_update_kernel_selector.h" #include "scatter_update/scatter_elements_update_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_nd_update.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_nd_update.cpp index 04fd75ef3c2..a05c8a3d8cd 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_nd_update.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_nd_update.cpp @@ -2,15 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "scatter_nd_update_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "scatter_nd_update_inst.h" #include "scatter_update/scatter_nd_update_kernel_selector.h" #include "scatter_update/scatter_nd_update_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_update.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_update.cpp index 44e5381a13d..26790d403ff 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_update.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_update.cpp @@ -2,15 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "scatter_update_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "scatter_update_inst.h" #include "scatter_update/scatter_update_kernel_selector.h" #include "scatter_update/scatter_update_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/select.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/select.cpp index 1440d4cdbea..684d43e28ec 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/select.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/select.cpp @@ -2,11 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "select_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "select_inst.h" #include "select/select_kernel_selector.h" #include "select/select_kernel_base.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/shape_of.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/shape_of.cpp index b5e5860763b..2fce15c96d5 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/shape_of.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/shape_of.cpp @@ -2,11 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "shape_of_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "intel_gpu/runtime/error_handler.hpp" -#include "kernel_selector_helper.h" + +#include "shape_of_inst.h" #include "shape_of/shape_of_kernel_selector.h" #include "shape_of/shape_of_kernel_ref.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/shuffle_channels.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/shuffle_channels.cpp index 64767a0f090..76bbe336f06 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/shuffle_channels.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/shuffle_channels.cpp @@ -2,15 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "shuffle_channels_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "shuffle_channels_inst.h" #include "shuffle_channels/shuffle_channels_kernel_selector.h" #include "shuffle_channels/shuffle_channels_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/slice.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/slice.cpp index 8a723b08e9f..898de26502a 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/slice.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/slice.cpp @@ -2,14 +2,13 @@ // SPDX-License-Identifier: Apache-2.0 // -#include -#include -#include -#include -#include -#include #include "primitive_base.hpp" -#include + +#include "slice_inst.h" +#include "data_inst.h" +#include "slice/slice_kernel_selector.h" +#include "slice/slice_kernel_ref.h" + #include #include @@ -30,7 +29,8 @@ std::vector extractIntegerData(const data_node& node, const stream } std::vector extractIntegerData(const data_node& node, const stream& stream) { - switch (node.get_output_layout().data_type) { + auto dt = node.get_output_layout().data_type; + switch (dt) { case data_types::u8: return extractIntegerData(node, stream); case data_types::i8: @@ -40,9 +40,7 @@ std::vector extractIntegerData(const data_node& node, const stream case data_types::i64: return extractIntegerData(node, stream); default: - CLDNN_ERROR_DATA_TYPES_MISMATCH(node.id(), "Slice parameter", - node.get_output_layout().data_type, "Any integral type", - data_types::i32, "Slice parameters should be of integral type."); + OPENVINO_ASSERT(false, "[GPU] Slice parameters should be of integral type for node ", node.id(), " while got ", dt); } return {}; } diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/softmax.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/softmax.cpp index 9f72fb5bfa3..d18dd4e83be 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/softmax.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/softmax.cpp @@ -2,13 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "softmax_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "softmax_inst.h" #include "softmax/softmax_kernel_selector.h" #include "softmax/softmax_kernel_base.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/space_to_batch.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/space_to_batch.cpp index 9cfb5c3a1da..0d23d3fa63d 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/space_to_batch.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/space_to_batch.cpp @@ -2,17 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "space_to_batch_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "space_to_batch_inst.h" #include "space_to_batch/space_to_batch_kernel_selector.h" #include "space_to_batch/space_to_batch_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" -#include "data_inst.h" -#include - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/space_to_depth.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/space_to_depth.cpp index 858c9e8c114..d60c67df2e8 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/space_to_depth.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/space_to_depth.cpp @@ -2,15 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "space_to_depth_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "space_to_depth_inst.h" #include "space_to_depth/space_to_depth_kernel_selector.h" #include "space_to_depth/space_to_depth_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/strided_slice.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/strided_slice.cpp index 05666905a19..989bc845f0b 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/strided_slice.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/strided_slice.cpp @@ -2,17 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "strided_slice_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "strided_slice_inst.h" +#include "data_inst.h" #include "strided_slice/strided_slice_kernel_ref.h" #include "strided_slice/strided_slice_kernel_selector.h" -#include "intel_gpu/runtime/error_handler.hpp" -#include "data_inst.h" -#include - -using namespace cldnn; namespace { template ::value>::type> diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/tile.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/tile.cpp index 43a24d21c1a..8c166bd1aaf 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/tile.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/tile.cpp @@ -2,15 +2,11 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "tile_inst.h" #include "primitive_base.hpp" -#include "impls/implementation_map.hpp" -#include "kernel_selector_helper.h" + +#include "tile_inst.h" #include "tile/tile_kernel_selector.h" #include "tile/tile_kernel_ref.h" -#include "intel_gpu/runtime/error_handler.hpp" - -using namespace cldnn; namespace cldnn { namespace ocl { diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/concatenation_onednn.cpp b/src/plugins/intel_gpu/src/graph/impls/onednn/concatenation_onednn.cpp index a516d617675..aa798b390f8 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/concatenation_onednn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/concatenation_onednn.cpp @@ -6,7 +6,7 @@ #include "eltwise_inst.h" #include "quantize_inst.h" #include "primitive_onednn_base.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "kernel_selector_common.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/convolution_onednn.cpp b/src/plugins/intel_gpu/src/graph/impls/onednn/convolution_onednn.cpp index f8d07513c57..2022d23972c 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/convolution_onednn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/convolution_onednn.cpp @@ -6,7 +6,7 @@ #include "eltwise_inst.h" #include "quantize_inst.h" #include "primitive_onednn_base.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "kernel_selector_common.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/deconvolution_onednn.cpp b/src/plugins/intel_gpu/src/graph/impls/onednn/deconvolution_onednn.cpp index c1227c57a3a..2d228d67344 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/deconvolution_onednn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/deconvolution_onednn.cpp @@ -6,7 +6,7 @@ #include "eltwise_inst.h" #include "quantize_inst.h" #include "primitive_onednn_base.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "kernel_selector_common.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/eltwise_onednn.cpp b/src/plugins/intel_gpu/src/graph/impls/onednn/eltwise_onednn.cpp index fcf2e405cd7..95d7f250964 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/eltwise_onednn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/eltwise_onednn.cpp @@ -4,7 +4,7 @@ #include "eltwise_inst.h" #include "primitive_onednn_base.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "kernel_selector_common.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.cpp b/src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.cpp index 62582f66693..b8b1fb70bd5 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.cpp @@ -4,7 +4,7 @@ #include "fully_connected_inst.h" #include "primitive_onednn_base.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "kernel_selector_common.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/gemm_onednn.cpp b/src/plugins/intel_gpu/src/graph/impls/onednn/gemm_onednn.cpp index 2b789fc34d4..2b74fb2ca00 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/gemm_onednn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/gemm_onednn.cpp @@ -4,7 +4,7 @@ #include "gemm_inst.h" #include "primitive_onednn_base.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "kernel_selector_common.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/pooling_onednn.cpp b/src/plugins/intel_gpu/src/graph/impls/onednn/pooling_onednn.cpp index 27cfbdbb29a..e1d592eecd7 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/pooling_onednn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/pooling_onednn.cpp @@ -4,7 +4,7 @@ #include "pooling_inst.h" #include "primitive_onednn_base.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "kernel_selector_common.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/primitive_onednn_base.h b/src/plugins/intel_gpu/src/graph/impls/onednn/primitive_onednn_base.h index 5e8c03dd0c6..30247fb7735 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/primitive_onednn_base.h +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/primitive_onednn_base.h @@ -8,7 +8,6 @@ #include "primitive_inst.h" #include "intel_gpu/graph/serialization/binary_buffer.hpp" -#include "intel_gpu/runtime/error_handler.hpp" #include "intel_gpu/runtime/memory.hpp" #include "to_string_utils.h" #include "register.hpp" diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/reduction_onednn.cpp b/src/plugins/intel_gpu/src/graph/impls/onednn/reduction_onednn.cpp index 4e26e945ef9..3e91c34c08b 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/reduction_onednn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/reduction_onednn.cpp @@ -4,7 +4,7 @@ #include "reduce_inst.h" #include "primitive_onednn_base.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "kernel_selector_common.h" #include "kernel_base.h" diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/reorder_onednn.cpp b/src/plugins/intel_gpu/src/graph/impls/onednn/reorder_onednn.cpp index a1a81a8955c..0bda1fa2f79 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/reorder_onednn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/reorder_onednn.cpp @@ -4,7 +4,7 @@ #include "reorder_inst.h" #include "primitive_onednn_base.h" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include "kernel_selector_common.h" diff --git a/src/plugins/intel_gpu/src/graph/include/assign_inst.h b/src/plugins/intel_gpu/src/graph/include/assign_inst.h index 90ca9c3bf8b..727c9b9a7c0 100644 --- a/src/plugins/intel_gpu/src/graph/include/assign_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/assign_inst.h @@ -6,7 +6,6 @@ #include "intel_gpu/primitives/assign.hpp" #include "primitive_inst.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { namespace memory_state { diff --git a/src/plugins/intel_gpu/src/graph/include/experimental_detectron_topk_rois_inst.h b/src/plugins/intel_gpu/src/graph/include/experimental_detectron_topk_rois_inst.h index c274521b256..c98a8ee3fbf 100644 --- a/src/plugins/intel_gpu/src/graph/include/experimental_detectron_topk_rois_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/experimental_detectron_topk_rois_inst.h @@ -6,7 +6,6 @@ #include "intel_gpu/primitives/experimental_detectron_topk_rois.hpp" #include "primitive_inst.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { diff --git a/src/plugins/intel_gpu/src/graph/include/eye_inst.h b/src/plugins/intel_gpu/src/graph/include/eye_inst.h index f380b6d96ab..a3e8a27b90a 100644 --- a/src/plugins/intel_gpu/src/graph/include/eye_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/eye_inst.h @@ -4,7 +4,6 @@ #pragma once #include -#include #include "primitive_inst.h" diff --git a/src/plugins/intel_gpu/src/graph/include/generic_layer.hpp b/src/plugins/intel_gpu/src/graph/include/generic_layer.hpp index 1d2d6e48bfc..67fe411e911 100644 --- a/src/plugins/intel_gpu/src/graph/include/generic_layer.hpp +++ b/src/plugins/intel_gpu/src/graph/include/generic_layer.hpp @@ -5,14 +5,14 @@ #pragma once #include "intel_gpu/primitives/primitive.hpp" #include "intel_gpu/runtime/memory.hpp" -#include "kernel_selector_helper.h" + +// TODO: Remove OCL impl dependency here or move to OCL folder +#include "impls/ocl/kernel_selector_helper.h" #include namespace cldnn { - - /// @brief Changes how data is ordered in memory. Value type is not changed & all information is preserved. /// @details Corresponding values are bitwise equal before/after reorder. /// Also merged with subtraction layer, which can subtract values while doing reordering. diff --git a/src/plugins/intel_gpu/src/graph/impls/implementation_map.hpp b/src/plugins/intel_gpu/src/graph/include/implementation_map.hpp similarity index 91% rename from src/plugins/intel_gpu/src/graph/impls/implementation_map.hpp rename to src/plugins/intel_gpu/src/graph/include/implementation_map.hpp index 431e1adbfd4..1282b0e05bb 100644 --- a/src/plugins/intel_gpu/src/graph/impls/implementation_map.hpp +++ b/src/plugins/intel_gpu/src/graph/include/implementation_map.hpp @@ -4,16 +4,15 @@ #pragma once +#include "intel_gpu/primitives/implementation_desc.hpp" +#include "kernel_impl_params.hpp" + #include #include -#include #include #include #include -#include "intel_gpu/primitives/implementation_desc.hpp" -#include "to_string_utils.h" -#include "kernel_selector_helper.h" namespace cldnn { @@ -30,25 +29,11 @@ public: } }; - -struct permute; -struct reorder; -struct custom_gpu_primitive; -struct generic_layer; -struct reshape; -struct data; -struct mutable_data; -struct input_layout; -struct prior_box; -struct loop; -struct shape_of; - struct primitive_impl; template struct typed_program_node; -template struct implementation_key { typedef std::tuple type; type operator()(const layout& proposed_layout) { @@ -56,16 +41,10 @@ struct implementation_key { } }; -namespace { -template -std::string get_key_name(const key_type &) { return std::string(""); } - -} // namespace - template class implementation_map { public: - using key_builder = implementation_key; + using key_builder = implementation_key; using key_type = typename key_builder::type; using factory_type = std::function(const typed_program_node&, const kernel_impl_params&)>; using map_type = singleton_map, std::pair, factory_type>>; @@ -87,7 +66,7 @@ public: } } OPENVINO_ASSERT(false, "[GPU] implementation_map for ", typeid(primitive_kind).name(), - " could not find any implementation to match key: ", get_key_name(key), + " could not find any implementation to match key: ", std::get<0>(key), "|", std::get<1>(key), ", impl_type: ", preferred_impl_type, ", shape_type: ", target_shape_type, ", node_id: ", impl_params.desc->id); } diff --git a/src/plugins/intel_gpu/src/graph/include/input_layout_inst.h b/src/plugins/intel_gpu/src/graph/include/input_layout_inst.h index dc1a3732a49..d8b8af440b5 100644 --- a/src/plugins/intel_gpu/src/graph/include/input_layout_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/input_layout_inst.h @@ -10,7 +10,6 @@ #include namespace cldnn { -struct memory; template <> struct typed_program_node : public typed_program_node_base { diff --git a/src/plugins/intel_gpu/src/graph/include/kernel_impl_params.hpp b/src/plugins/intel_gpu/src/graph/include/kernel_impl_params.hpp new file mode 100644 index 00000000000..079441d8afa --- /dev/null +++ b/src/plugins/intel_gpu/src/graph/include/kernel_impl_params.hpp @@ -0,0 +1,124 @@ +// Copyright (C) 2023 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#pragma once + +#include "intel_gpu/graph/serialization/binary_buffer.hpp" +#include "intel_gpu/runtime/engine.hpp" +#include "intel_gpu/runtime/memory.hpp" +#include "intel_gpu/runtime/utils.hpp" +#include "intel_gpu/runtime/tensor.hpp" +#include "intel_gpu/primitives/eltwise.hpp" +#include "intel_gpu/primitives/quantize.hpp" +#include "intel_gpu/primitives/activation.hpp" +#include "intel_gpu/primitives/primitive.hpp" + +#include "tensor_type.h" +#include "fused_primitive_desc.h" + +#include +#include +#include +#include + +namespace cldnn { + +struct program; + + +struct kernel_impl_params { + bool has_runtime_layouts = false; + const program *prog; + std::shared_ptr desc; + size_t unique_id; + std::vector input_layouts; + std::vector output_layouts; + std::vector input_offsets; + std::vector fused_desc; +#ifdef ENABLE_ONEDNN_FOR_GPU + std::vector fused_desc_onednn; +#endif // ENABLE_ONEDNN_FOR_GPU + + optional_layout weights_layout = optional_layout(); + + optional_layout bias_layout = optional_layout(); + optional_layout weights_zero_points_layout = optional_layout(); + optional_layout activations_zero_points_layout = optional_layout(); + optional_layout compensation_layout = optional_layout(); + + std::map memory_deps = {}; + size_t primary_input_idx = 0; + + memory::ptr reordered_weights = nullptr; + + kernel_impl_params() {} + + kernel_impl_params(program& _prog, + std::shared_ptr _desc, + size_t _uid, + const std::vector& _in_layouts, + const std::vector& _out_layouts, + const std::vector& _fused_descs) + : has_runtime_layouts(true) + , prog(&_prog) + , desc(_desc) + , unique_id(_uid) + , input_layouts(_in_layouts) + , output_layouts(_out_layouts) + , fused_desc(_fused_descs) + , primary_input_idx(0) { + } + + layout get_input_layout(size_t idx = 0) const { + OPENVINO_ASSERT(input_layouts.size() > idx, + "The size of input layouts must be greater than the requested index: ", + "Requested index is ", idx, ", ", + "but the size of input layouts is ", input_layouts.size()); + return input_layouts[idx]; + } + + layout get_non_padded_input_layout(size_t idx = 0) const { + auto input_layout = get_input_layout(idx); + auto result = layout({input_layout.get_partial_shape(), input_layout.data_type, input_layout.format}); + return result; + } + + layout get_output_layout(size_t idx = 0) const { + OPENVINO_ASSERT(output_layouts.size() > idx, + "The size of output layouts must be greater than the requested index: ", + "Requested index is ", idx, ",", + "but the size of output layouts is ", output_layouts.size()); + return output_layouts[idx]; + } + + bool has_fused_primitives() const { return !fused_desc.empty(); } + + layout get_fused_output_layout() const { + if (fused_desc.empty()) + return layout(data_types::f32, format::bfyx, tensor()); + return fused_desc.back().output_layout; + } + + bool is_dynamic() const { + for (auto i : input_layouts) + if (i.is_dynamic()) + return true; + for (auto i : output_layouts) + if (i.is_dynamic()) + return true; + return false; + } + + template + std::shared_ptr typed_desc() const { return std::static_pointer_cast(desc); } + + void save(BinaryOutputBuffer& ob) const; + void load(BinaryInputBuffer& ib); + const program& get_program() const { + OPENVINO_ASSERT(prog != nullptr, "[GPU] Program pointer in kernel_impl_params in not initialized"); + return *prog; + } +}; + +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/graph/include/layout_optimizer.h b/src/plugins/intel_gpu/src/graph/include/layout_optimizer.h index 21fd41a59c8..f07b03c624c 100644 --- a/src/plugins/intel_gpu/src/graph/include/layout_optimizer.h +++ b/src/plugins/intel_gpu/src/graph/include/layout_optimizer.h @@ -20,8 +20,8 @@ #include "non_max_suppression_inst.h" #include "region_yolo_inst.h" -#include "kernel_selector_common.h" -#include "kernel_selector_helper.h" +// TODO: add generic interface for weights_reorder_params and get rid of this dependency +#include "impls/ocl/kernel_selector_helper.h" #include #include diff --git a/src/plugins/intel_gpu/src/graph/include/loop_inst.h b/src/plugins/intel_gpu/src/graph/include/loop_inst.h index f05ad3ae8bb..15f18f4b6c7 100644 --- a/src/plugins/intel_gpu/src/graph/include/loop_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/loop_inst.h @@ -10,7 +10,6 @@ #include "intel_gpu/primitives/input_layout.hpp" #include "intel_gpu/primitives/eltwise.hpp" #include "intel_gpu/runtime/memory.hpp" -#include "intel_gpu/runtime/error_handler.hpp" #include "primitive_inst.h" #include @@ -221,11 +220,8 @@ public: } else { const auto& body_input_prim = body.at(current_iteration_id); const auto input_layout_prim = std::dynamic_pointer_cast(body_input_prim); - if (!input_layout_prim) { - CLDNN_ERROR_MESSAGE(this->id(), "current_iteration primitive should be cldnn::input_layout"); - } else { - input_layout_prim->change_layout(body_input_layout); - } + OPENVINO_ASSERT(input_layout_prim, "[GPU] current_iteration primitive should be cldnn::input_layout in node", this->id()); + input_layout_prim->change_layout(body_input_layout); } // add incremental data: 1 @@ -277,9 +273,7 @@ public: } // setup internal output - if (output_primitive_maps.empty()) { - CLDNN_ERROR_MESSAGE(this->id(), "output primitive map should have at least 1 mapping"); - } + OPENVINO_ASSERT(!output_primitive_maps.empty(), "[GPU] Output primitive map should have at least 1 mapping in primitive ", this->id()); std::set output_names; output_names.insert(output_primitive_maps.front().internal_id); @@ -299,8 +293,8 @@ public: // input primitive map because its initial value is always // zero and the value will be set in execute_impl() if (back_edge.to != get_current_iteration_id() && input_map == input_primitive_maps.end()) { - std::string msg = "No primitive mapping for backedge (internal_id: " + back_edge.to + ')'; - CLDNN_ERROR_MESSAGE(this->id(), msg.c_str()); + std::string msg = "[GPU] No primitive mapping for backedge (internal_id: " + back_edge.to + ") for primitive " + this->id(); + OPENVINO_ASSERT(false, msg.c_str()); } output_names.insert(back_edge.from); @@ -536,9 +530,7 @@ public: void update_mapped_memory(); void set_output_memory(memory::ptr mem, bool check = true, size_t idx = 0) override; const backedge_memory_mapping& get_current_iteration_backedge_mapping() const { - if (!node->is_current_iteration_used()) { - CLDNN_ERROR_MESSAGE(node->id(), "no backedge mapping for current_iteration"); - } + OPENVINO_ASSERT(node->is_current_iteration_used(), "[GPU] No backedge mapping for current_iteration for primitive ", node->id()); return backedge_memory_mappings.at(current_iteratoin_backedge_mapping_idx); } void save(BinaryOutputBuffer& ob) const override; diff --git a/src/plugins/intel_gpu/src/graph/include/lstm_dynamic_input_inst.h b/src/plugins/intel_gpu/src/graph/include/lstm_dynamic_input_inst.h index 329cd2d658b..0dece56e7d2 100644 --- a/src/plugins/intel_gpu/src/graph/include/lstm_dynamic_input_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/lstm_dynamic_input_inst.h @@ -5,6 +5,7 @@ #pragma once #include "intel_gpu/primitives/lstm_dynamic_input.hpp" #include "primitive_inst.h" +#include "intel_gpu/runtime/error_handler.hpp" #include #include diff --git a/src/plugins/intel_gpu/src/graph/include/lstm_dynamic_inst.h b/src/plugins/intel_gpu/src/graph/include/lstm_dynamic_inst.h index c334bedf3d7..011c226cda9 100644 --- a/src/plugins/intel_gpu/src/graph/include/lstm_dynamic_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/lstm_dynamic_inst.h @@ -5,6 +5,7 @@ #pragma once #include "intel_gpu/primitives/lstm_dynamic.hpp" #include "primitive_inst.h" +#include "intel_gpu/runtime/error_handler.hpp" #include #include diff --git a/src/plugins/intel_gpu/src/graph/include/non_zero_inst.h b/src/plugins/intel_gpu/src/graph/include/non_zero_inst.h index ca69fd3e003..a8248e15931 100644 --- a/src/plugins/intel_gpu/src/graph/include/non_zero_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/non_zero_inst.h @@ -5,7 +5,6 @@ #pragma once #include "intel_gpu/primitives/non_zero.hpp" #include "primitive_inst.h" -#include "intel_gpu/runtime/error_handler.hpp" #include #include diff --git a/src/plugins/intel_gpu/src/graph/include/primitive_inst.h b/src/plugins/intel_gpu/src/graph/include/primitive_inst.h index c21ad7d3906..50ab0b037ce 100644 --- a/src/plugins/intel_gpu/src/graph/include/primitive_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/primitive_inst.h @@ -5,11 +5,9 @@ #pragma once #include "intel_gpu/primitives/primitive.hpp" #include "intel_gpu/primitives/concatenation.hpp" -#include "intel_gpu/runtime/error_handler.hpp" #include "intel_gpu/runtime/event.hpp" #include "intel_gpu/runtime/memory.hpp" #include "intel_gpu/graph/network.hpp" -#include "kernel_selector_helper.h" #include "meta_utils.h" #include "program_node.h" #include "primitive_type.h" @@ -22,6 +20,9 @@ #include "intel_gpu/graph/serialization/vector_serializer.hpp" #include "runtime/kernels_cache.hpp" +// TODO: add generic interface for weights_reorder_params and get rid of this dependency +#include "impls/ocl/kernel_selector_helper.h" + #include #include #include @@ -43,6 +44,8 @@ struct primitive_impl { primitive_impl() = default; explicit primitive_impl(const kernel_selector::weights_reorder_params& params, std::string kernel_name = "", bool is_dynamic = false) : _weights_reorder_params(params), _kernel_name(kernel_name), _is_dynamic(is_dynamic) {} + explicit primitive_impl(std::string kernel_name, bool is_dynamic = false) : + primitive_impl(kernel_selector::weights_reorder_params{}, kernel_name, is_dynamic) {} virtual ~primitive_impl() = default; virtual std::vector get_internal_buffer_layouts() const = 0; diff --git a/src/plugins/intel_gpu/src/graph/include/primitive_type.h b/src/plugins/intel_gpu/src/graph/include/primitive_type.h index 449d36e59e9..475da762b5a 100644 --- a/src/plugins/intel_gpu/src/graph/include/primitive_type.h +++ b/src/plugins/intel_gpu/src/graph/include/primitive_type.h @@ -6,7 +6,7 @@ #include "intel_gpu/runtime/layout.hpp" #include "intel_gpu/runtime/memory.hpp" -#include "kernel_selector_helper.h" +#include "kernel_impl_params.hpp" #include #include diff --git a/src/plugins/intel_gpu/src/graph/include/primitive_type_base.h b/src/plugins/intel_gpu/src/graph/include/primitive_type_base.h index f7504e102b2..7e859f65ac5 100644 --- a/src/plugins/intel_gpu/src/graph/include/primitive_type_base.h +++ b/src/plugins/intel_gpu/src/graph/include/primitive_type_base.h @@ -13,7 +13,7 @@ #include "program_node.h" #include "primitive_inst.h" #include "intel_gpu/graph/network.hpp" -#include "impls/implementation_map.hpp" +#include "implementation_map.hpp" #include #include @@ -42,11 +42,21 @@ struct primitive_type_base : primitive_type { } std::unique_ptr choose_impl(const cldnn::program_node& node, const kernel_impl_params& runtime_params) const override { - OPENVINO_ASSERT(node.type() == this, "[GPU] primitive_type_base::choose_impl: primitive type mismatch"); - auto factory = implementation_map::get(runtime_params, node.get_preferred_impl_type(), get_shape_type(runtime_params)); - auto impl = factory(node, runtime_params); - impl->set_dynamic(get_shape_type(runtime_params) == shape_types::dynamic_shape); - return impl; + try { + OPENVINO_ASSERT(node.type() == this, "[GPU] primitive_type_base::choose_impl: primitive type mismatch"); + auto factory = implementation_map::get(runtime_params, node.get_preferred_impl_type(), get_shape_type(runtime_params)); + auto impl = factory(node, runtime_params); + impl->set_dynamic(get_shape_type(runtime_params) == shape_types::dynamic_shape); + return impl; + } catch (std::exception& e) { + std::stringstream ss; + const auto& p = node.get_primitive(); + ov::write_all_to_stream(ss, "[GPU] Can't choose implementation for ", node.id(), " node (type=", p->type_string(), ")\n", + "[GPU] Original name: ", p->origin_op_name, "\n" + "[GPU] Original type: ", p->origin_op_type_name, "\n" + "[GPU] Reason: ", e.what()); + throw ov::Exception(ss.str()); + } } bool does_an_implementation_exist(const cldnn::program_node& node) const override { diff --git a/src/plugins/intel_gpu/src/graph/include/program_node.h b/src/plugins/intel_gpu/src/graph/include/program_node.h index 8042bebb57a..0d3e3427466 100644 --- a/src/plugins/intel_gpu/src/graph/include/program_node.h +++ b/src/plugins/intel_gpu/src/graph/include/program_node.h @@ -9,8 +9,8 @@ #include "intel_gpu/primitives/implementation_desc.hpp" #include "intel_gpu/graph/program.hpp" -#include "kernel_selector_helper.h" #include "fused_primitive_desc.h" +#include "kernel_impl_params.hpp" #include "meta_utils.h" #include @@ -21,6 +21,11 @@ #include #include +// TODO: Remove forward declarations for kernel_selector once fused ops descriptors don't depend on OCL stuff +namespace kernel_selector { +struct fuse_params; +} + namespace cldnn { struct program; diff --git a/src/plugins/intel_gpu/src/graph/include/random_uniform_inst.h b/src/plugins/intel_gpu/src/graph/include/random_uniform_inst.h index 868fb3265e4..4768a702d08 100644 --- a/src/plugins/intel_gpu/src/graph/include/random_uniform_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/random_uniform_inst.h @@ -6,7 +6,6 @@ #include "intel_gpu/primitives/random_uniform.hpp" #include "primitive_inst.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { diff --git a/src/plugins/intel_gpu/src/graph/include/read_value_inst.h b/src/plugins/intel_gpu/src/graph/include/read_value_inst.h index cf34f1da5e2..9e5c426d4a6 100644 --- a/src/plugins/intel_gpu/src/graph/include/read_value_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/read_value_inst.h @@ -7,7 +7,6 @@ #include "assign_inst.h" #include "intel_gpu/primitives/read_value.hpp" #include "primitive_inst.h" -#include "intel_gpu/runtime/error_handler.hpp" namespace cldnn { diff --git a/src/plugins/intel_gpu/src/graph/include/reshape_inst.h b/src/plugins/intel_gpu/src/graph/include/reshape_inst.h index beb3b743b8b..1b55303cfaf 100644 --- a/src/plugins/intel_gpu/src/graph/include/reshape_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/reshape_inst.h @@ -5,7 +5,6 @@ #pragma once #include "intel_gpu/primitives/reshape.hpp" #include "primitive_inst.h" -#include "intel_gpu/runtime/error_handler.hpp" #include #include @@ -22,10 +21,7 @@ struct typed_program_node : public typed_program_node_base { public: using parent::parent; - program_node& input() const { - CLDNN_ERROR_LESS_THAN(id(), "the number of dependencies", dependencies.size(), "1", 1, "ERROR: the node has no input"); - return get_dependency(0); - } + program_node& input() const { return get_dependency(0); } bool is_in_place() const { if (this->is_output() || this->has_fused_primitives()) diff --git a/src/plugins/intel_gpu/src/graph/include/roi_align_inst.h b/src/plugins/intel_gpu/src/graph/include/roi_align_inst.h index 1503a1ba3dd..5a9801d5aec 100644 --- a/src/plugins/intel_gpu/src/graph/include/roi_align_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/roi_align_inst.h @@ -5,7 +5,6 @@ #include #include "primitive_inst.h" -#include namespace cldnn { diff --git a/src/plugins/intel_gpu/src/graph/include/slice_inst.h b/src/plugins/intel_gpu/src/graph/include/slice_inst.h index 20dffc89961..fe11a7058ad 100644 --- a/src/plugins/intel_gpu/src/graph/include/slice_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/slice_inst.h @@ -5,7 +5,6 @@ #pragma once #include #include "primitive_inst.h" -#include namespace cldnn { diff --git a/src/plugins/intel_gpu/src/graph/include/sliding_window_utils.hpp b/src/plugins/intel_gpu/src/graph/include/sliding_window_utils.hpp index c8b2df8862b..e39007e7060 100644 --- a/src/plugins/intel_gpu/src/graph/include/sliding_window_utils.hpp +++ b/src/plugins/intel_gpu/src/graph/include/sliding_window_utils.hpp @@ -75,7 +75,7 @@ tensor calc_sliding_window_output_range(const tensor&, const tensor&, bool, const tensor::value_type&) { - static_assert(meta::always_false>::value, + static_assert(cldnn::meta::always_false>::value, "Sliding window output range calculation mode is not supported. Please implement specialization " "for new swor_mode."); diff --git a/src/plugins/intel_gpu/src/graph/include/sliding_window_utils_legacy.h b/src/plugins/intel_gpu/src/graph/include/sliding_window_utils_legacy.h deleted file mode 100644 index 9dc3025e333..00000000000 --- a/src/plugins/intel_gpu/src/graph/include/sliding_window_utils_legacy.h +++ /dev/null @@ -1,389 +0,0 @@ -// Copyright (C) 2018-2023 Intel Corporation -// SPDX-License-Identifier: Apache-2.0 -// - -#pragma once - -#include "intel_gpu/runtime/layout.hpp" -#include "intel_gpu/runtime/tensor.hpp" - -#include "meta_utils.h" - -#include -#include -#include - -namespace cldnn { - -/// @brief Sliding window output range computation mode. -enum class swor_mode { - // Single modes: - all, ///< Range is computed in the way that each sliding window in range is fully contained inside - ///< (optionally upper-padded by pad) input data. - exceed_once, ///< Range is computed in the way that each except at most one sliding window in range is fully - ///< contained inside (optionally upper-padded by pad) input data. The last window may partially - ///< exceed (optionally upper-padded by pad) input data range. - any, ///< Range is computed in the way that each sliding window in range is fully or at least partially - ///< contained inside (optionally upper-padded by pad) input data. - // Mixed modes: - exceed_once_data, ///< Range is computed in the way that each except at most one sliding window in range is fully - ///< contained inside (optionally upper-padded by pad) input data. The last window may - ///< partially exceed (non-upper-padded) input data range. - ///< This mode is effectievely minimum of combination of @c swor_mode::exceed_once mode - ///< and @c swor_mode::any mode (with always @c sym_pad = false). - max ///< Maximum of all single modes with all cominations of @c sym_pad. -}; - -/// @brief Calculates output range (size) for sliding window moving on input data range specified by @p input_size. -/// -/// @param input_size Range/Size of input data (non-padded or treated as valid). Only spatial coordinates are -/// considered. -/// @param size Size of sliding window. Only spatial coordinates are considered. -/// @param pad pad/Padding of sliding window in input. Only spatial coordinates are considered. Padding/pad -/// is applied from both sides of input data: negative value extends/pads data, positive - crops it. -/// @param stride Horizontal/Vertical stride of sliding in input data. -/// @param dilation Horizontal/Vertical dilation of sliding window on input data. -/// @param sym_pad Treat pad as applied on input symmetrically (from both sides). If @c false, the @p pad -/// is applied only from left/upper side. -/// @param degen_val If values from calculation are in allowed range, but calculated output size is invalid, -/// the @p degen_val is returned. Any non-positive value is considered degenerated and will be -/// switched to value passed in this parameter. -/// @return Output range (size) of sliding window. Only spatial dimensions are valid (rest is 0). -template -tensor calc_sliding_window_output_range(const tensor& input_size, - const tensor& size, - const tensor& pad, - const tensor& stride, - const tensor& dilation = {1, 1, 1, 1}, - bool sym_pad = true, - const tensor::value_type& degen_val = 0); - -/// @brief Fall-back implementation. -template -tensor calc_sliding_window_output_range(const tensor&, - const tensor&, - const tensor&, - const tensor&, - const tensor&, - bool, - const tensor::value_type&) { - static_assert(meta::always_false>::value, - "Sliding window output range calculation mode is not supported. Please implement specialization " - "for new swor_mode."); - - return tensor(); -} - -template <> -inline tensor calc_sliding_window_output_range(const tensor& input_size, - const tensor& size, - const tensor& pad, - const tensor& stride, - const tensor& dilation, - bool sym_pad, - const tensor::value_type& degen_val) { - if (input_size.spatial[0] <= 0 || input_size.spatial[1] <= 0 || input_size.spatial[2] <= 0) - throw std::invalid_argument("Input data spatial sizes must be positive (>= 1)."); - if (size.spatial[0] <= 0 || size.spatial[1] <= 0 || size.spatial[2] <= 0) - throw std::invalid_argument("Sliding window spatial sizes must be positive (>= 1)."); - if (stride.spatial[0] <= 0 || stride.spatial[1] <= 0 || stride.spatial[2] <= 0) - throw std::invalid_argument("Sliding window h/v strides must be positive (>= 1)."); - if (dilation.spatial[0] <= 0 || dilation.spatial[1] <= 0 || dilation.spatial[2] <= 0) - throw std::invalid_argument("Sliding window h/v input dialations must be positive (>= 1)."); - - auto off_factor = sym_pad ? -2 : -1; - tensor wnd_ext_size{0, - 0, - (size.spatial[0] - 1) * dilation.spatial[0] + 1, - (size.spatial[1] - 1) * dilation.spatial[1] + 1, - (size.spatial[2] - 1) * dilation.spatial[2] + 1}; - - // wes = (size - 1) * dilation + 1 - // lpos(i) = -pad + i * stride + wes - 1, for i = 0, 1, ... - // - // output_range = max {i | lpos(i) < input_size + pad} + 1, if sym_pad is true - // output_range = max {i | lpos(i) < input_size} + 1, if sym_pad is false - auto output_range_x = static_cast( - off_factor * pad.spatial[0] + wnd_ext_size.spatial[0] <= input_size.spatial[0] - ? (input_size.spatial[0] - off_factor * pad.spatial[0] - wnd_ext_size.spatial[0]) / stride.spatial[0] + 1 - : degen_val); - auto output_range_y = static_cast( - off_factor * pad.spatial[1] + wnd_ext_size.spatial[1] <= input_size.spatial[1] - ? (input_size.spatial[1] - off_factor * pad.spatial[1] - wnd_ext_size.spatial[1]) / stride.spatial[1] + 1 - : degen_val); - auto output_range_z = static_cast( - off_factor * pad.spatial[2] + wnd_ext_size.spatial[2] <= input_size.spatial[2] - ? (input_size.spatial[2] - off_factor * pad.spatial[2] - wnd_ext_size.spatial[2]) / stride.spatial[2] + 1 - : degen_val); - - return {0, 0, output_range_x, output_range_y, output_range_z}; -} - -template <> -inline tensor calc_sliding_window_output_range(const tensor& input_size, - const tensor& size, - const tensor& pad, - const tensor& stride, - const tensor& dilation, - bool sym_pad, - const tensor::value_type& degen_val) { - if (input_size.spatial[0] <= 0 || input_size.spatial[1] <= 0 || input_size.spatial[2] <= 0) - throw std::invalid_argument("Input data spatial sizes must be positive (>= 1)."); - if (size.spatial[0] <= 0 || size.spatial[1] <= 0 || size.spatial[2] <= 0) - throw std::invalid_argument("Sliding window spatial sizes must be positive (>= 1)."); - if (stride.spatial[0] <= 0 || stride.spatial[1] <= 0 || stride.spatial[2] <= 0) - throw std::invalid_argument("Sliding window h/v strides must be positive (>= 1)."); - if (dilation.spatial[0] <= 0 || dilation.spatial[1] <= 0 || dilation.spatial[2] <= 0) - throw std::invalid_argument("Sliding window h/v input dialations must be positive (>= 1)."); - - auto off_factor = sym_pad ? -2 : -1; - tensor wnd_ext_size{0, - 0, - (size.spatial[0] - 1) * dilation.spatial[0] + 1, - (size.spatial[1] - 1) * dilation.spatial[1] + 1, - (size.spatial[2] - 1) * dilation.spatial[2] + 1}; - - tensor extend = tensor::max(wnd_ext_size, stride); - - // wes = (size - 1) * dilation + 1 - // fpos(i) = -pad + i * stride, for i = 0, 1, ... - // lpos(i) = -pad + i * stride + wes - 1, for i = 0, 1, ... - // - // output_range = max {i | lpos(i) < input_size + pad - 1 and fpos(i + 1) < input_size + pad} + 2, if - // sym_pad is true output_range = max {i | lpos(i) < input_size - 1 and fpos(i + 1) < input_size} + 2, - // if sym_pad is false - auto output_range_x = static_cast( - off_factor * pad.spatial[0] + extend.spatial[0] <= input_size.spatial[0] + stride.spatial[0] - 1 - ? (input_size.spatial[0] - off_factor * pad.spatial[0] - extend.spatial[0] + stride.spatial[0] - 1) / - stride.spatial[0] + - 1 - : degen_val); - auto output_range_y = static_cast( - off_factor * pad.spatial[1] + extend.spatial[1] <= input_size.spatial[1] + stride.spatial[1] - 1 - ? (input_size.spatial[1] - off_factor * pad.spatial[1] - extend.spatial[1] + stride.spatial[1] - 1) / - stride.spatial[1] + - 1 - : degen_val); - auto output_range_z = static_cast( - off_factor * pad.spatial[2] + extend.spatial[2] <= input_size.spatial[2] + stride.spatial[2] - 1 - ? (input_size.spatial[2] - off_factor * pad.spatial[2] - extend.spatial[2] + stride.spatial[2] - 1) / - stride.spatial[2] + - 1 - : degen_val); - - return {0, 0, output_range_x, output_range_y, output_range_z}; -} - -template <> -inline tensor calc_sliding_window_output_range(const tensor& input_size, - const tensor& size, - const tensor& pad, - const tensor& stride, - const tensor& dilation, - bool sym_pad, - const tensor::value_type& degen_val) { - if (input_size.spatial[0] <= 0 || input_size.spatial[1] <= 0 || input_size.spatial[2] <= 0) - throw std::invalid_argument("Input data spatial sizes must be positive (>= 1)."); - if (size.spatial[0] <= 0 || size.spatial[1] <= 0 || size.spatial[2] <= 0) - throw std::invalid_argument("Sliding window spatial sizes must be positive (>= 1)."); - if (stride.spatial[0] <= 0 || stride.spatial[1] <= 0 || stride.spatial[2] <= 0) - throw std::invalid_argument("Sliding window h/v strides must be positive (>= 1)."); - if (dilation.spatial[0] <= 0 || dilation.spatial[1] <= 0 || dilation.spatial[2] <= 0) - throw std::invalid_argument("Sliding window h/v input dialations must be positive (>= 1)."); - - auto off_factor = sym_pad ? -2 : -1; - - // fpos(i) = -pad + i * stride, for i = 0, 1, ... - // - // output_range = max {i | fpos(i) < input_size + pad} + 1, if sym_pad is true - // output_range = max {i | fpos(i) < input_size} + 1, if sym_pad is false - auto output_range_x = static_cast( - off_factor * pad.spatial[0] <= input_size.spatial[0] - 1 - ? (input_size.spatial[0] - off_factor * pad.spatial[0] - 1) / stride.spatial[0] + 1 - : degen_val); - auto output_range_y = static_cast( - off_factor * pad.spatial[1] <= input_size.spatial[1] - 1 - ? (input_size.spatial[1] - off_factor * pad.spatial[1] - 1) / stride.spatial[1] + 1 - : degen_val); - auto output_range_z = static_cast( - off_factor * pad.spatial[2] <= input_size.spatial[2] - 1 - ? (input_size.spatial[2] - off_factor * pad.spatial[2] - 1) / stride.spatial[2] + 1 - : degen_val); - - return {0, 0, output_range_x, output_range_y, output_range_z}; -} - -template <> -inline tensor calc_sliding_window_output_range(const tensor& input_size, - const tensor& size, - const tensor& pad, - const tensor& stride, - const tensor& dilation, - bool sym_pad, - const tensor::value_type& degen_val) { - auto output_range_exceed_once = calc_sliding_window_output_range(input_size, - size, - pad, - stride, - dilation, - sym_pad, - degen_val); - auto output_range_exceed_any_data = - calc_sliding_window_output_range(input_size, size, pad, stride, dilation, false, degen_val); - - return tensor::min(output_range_exceed_once, output_range_exceed_any_data); -} - -template <> -inline tensor calc_sliding_window_output_range(const tensor& input_size, - const tensor& size, - const tensor& pad, - const tensor& stride, - const tensor& dilation, - bool, - const tensor::value_type& degen_val) { - auto output_range_all_sym = - calc_sliding_window_output_range(input_size, size, pad, stride, dilation, true, degen_val); - auto output_range_all_asym = - calc_sliding_window_output_range(input_size, size, pad, stride, dilation, false, degen_val); - - auto output_range_exceed_once_sym = calc_sliding_window_output_range(input_size, - size, - pad, - stride, - dilation, - true, - degen_val); - auto output_range_exceed_once_asym = calc_sliding_window_output_range(input_size, - size, - pad, - stride, - dilation, - false, - degen_val); - - auto output_range_any_sym = - calc_sliding_window_output_range(input_size, size, pad, stride, dilation, true, degen_val); - auto output_range_any_asym = - calc_sliding_window_output_range(input_size, size, pad, stride, dilation, false, degen_val); - - return tensor::max(tensor::max(tensor::max(output_range_all_sym, output_range_all_asym), - tensor::max(output_range_exceed_once_sym, output_range_exceed_once_asym)), - tensor::max(output_range_any_sym, output_range_any_asym)); -} - -/// @brief Calculates minumum needed input range (size) for sliding window to get at least specified @p output_size. -/// -/// @param output_size Range/Size of output data (non-padded or treated as valid). Only spatial coordinates are -/// considered. -/// @param size Size of sliding window. Only spatial coordinates are considered. -/// @param pad pad/Padding of sliding window in input. Only spatial coordinates are considered. Padding/pad -/// is applied from both sides of input data: negative value extends/pads data, positive - crops it. -/// @param stride Horizontal/Vertical stride of sliding in input data. -/// @param dilation Horizontal/Vertical dilation of sliding window on input data. -/// @param sym_pad Treat pad as applied on input symmetrically (from both sides). If @c false, the @p pad -/// is applied only from left/upper side. -/// @param degen_val If values from calculation are in allowed range, but calculated output size is invalid, -/// the @p degen_val is returned. Any non-positive value is considered degenerated and will be -/// switched to value passed in this parameter. -/// @return Input range (size) for sliding window to get equal or greater @p output_size. -inline tensor calc_sliding_window_needed_input_range(const tensor& output_size, - const tensor& size, - const tensor& pad, - const tensor& stride, - const tensor& dilation = {1, 1, 1, 1}, - bool sym_pad = true, - const tensor::value_type& degen_val = 0) { - if (output_size.spatial[0] <= 0 || output_size.spatial[1] <= 0 || output_size.spatial[2] <= 0) - throw std::invalid_argument("Output data spatial sizes must be positive (>= 1)."); - if (size.spatial[0] <= 0 || size.spatial[1] <= 0 || size.spatial[2] <= 0) - throw std::invalid_argument("Sliding window spatial sizes must be positive (>= 1)."); - if (stride.spatial[0] <= 0 || stride.spatial[1] <= 0 || stride.spatial[2] <= 0) - throw std::invalid_argument("Sliding window h/v strides must be positive (>= 1)."); - if (dilation.spatial[0] <= 0 || dilation.spatial[1] <= 0 || dilation.spatial[2] <= 0) - throw std::invalid_argument("Sliding window h/v input dialations must be positive (>= 1)."); - - auto off_factor = sym_pad ? -2 : -1; - tensor wnd_ext_size{0, - 0, - (size.spatial[0] - 1) * dilation.spatial[0] + 1, - (size.spatial[1] - 1) * dilation.spatial[1] + 1, - (size.spatial[2] - 1) * dilation.spatial[2] + 1}; - - auto output_range_x = - off_factor * pad.spatial[0] + (output_size.spatial[0] - 1) * stride.spatial[0] + wnd_ext_size.spatial[0]; - auto output_range_y = - off_factor * pad.spatial[1] + (output_size.spatial[1] - 1) * stride.spatial[1] + wnd_ext_size.spatial[1]; - auto output_range_z = - off_factor * pad.spatial[2] + (output_size.spatial[2] - 1) * stride.spatial[2] + wnd_ext_size.spatial[2]; - - if (output_range_x <= 0) - output_range_x = degen_val; - if (output_range_y <= 0) - output_range_y = degen_val; - if (output_range_z <= 0) - output_range_z = degen_val; - - return {0, 0, output_range_x, output_range_y, output_range_z}; -} - -/// @brief Calculates safe needed input upper padding for sliding window to be able to compute at least -/// specified @p output_size. -/// -/// @param output_size Range/Size of output data (non-padded or treated as valid). Only spatial coordinates are -/// considered. -/// @param size Size of sliding window. Only spatial coordinates are considered. -/// @param pad Padding of sliding window in input. Only spatial coordinates are considered. Padding/pad -/// is applied from both sides of input data: negative value extends/pads data, positive - crops it. -/// @param stride Horizontal/Vertical stride of sliding in input data. -/// @param dilation Horizontal/Vertical dilation of sliding window on input data. -/// @param inverse Indicate that inverse calculation of needed range should take place (estimation of needed -/// ouput size when input size is specified). Used in deconvolution (when we switch input calculation -/// with output calculation). -/// @param degen_val If values from calculation are in allowed range, but calculated output size is invalid, -/// the @p degen_val is returned. Any non-positive value is considered degenerated and will be -/// switched to value passed in this parameter. -/// @return Input upper padding for sliding window to get equal or greater @p output_size. The padding takes into -/// consideration actual value of padding (always extends it) and only works on spatial coordinates of upper -/// padding (rest of padding values are not changed). -inline padding calc_sliding_window_needed_input_padding(const layout& actual_input_layout, - const tensor& output_size, - const tensor& size, - const tensor& pad, - const tensor& stride, - const tensor& dilation = {1, 1, 1, 1}, - bool inverse = false, - const tensor::value_type& degen_val = 0) { - tensor needed_size; - if (inverse) { - needed_size = calc_sliding_window_output_range(output_size, - size, - pad, - stride, - dilation, - false /* not important */, - degen_val); - } else { - auto needed_size_sym = - calc_sliding_window_needed_input_range(output_size, size, pad, stride, dilation, true, degen_val); - auto needed_size_asym = - calc_sliding_window_needed_input_range(output_size, size, pad, stride, dilation, false, degen_val); - - needed_size = tensor::max(needed_size_sym, needed_size_asym); - } - - const auto& actual_data_size = actual_input_layout.get_tensor(); - const auto& actual_lpad = actual_input_layout.data_padding.lower_size(); - const auto& actual_upad = actual_input_layout.data_padding.upper_size(); - - auto needed_upad = tensor::max(needed_size.sub(actual_data_size), actual_upad); - - return padding(actual_lpad.sizes(), - {actual_upad.batch[0], - actual_upad.feature[0], - needed_upad.spatial[0], - needed_upad.spatial[1], - needed_upad.spatial[2]}); -} - -} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/graph/include/strided_slice_inst.h b/src/plugins/intel_gpu/src/graph/include/strided_slice_inst.h index b685917d7be..dfcd6bb36aa 100644 --- a/src/plugins/intel_gpu/src/graph/include/strided_slice_inst.h +++ b/src/plugins/intel_gpu/src/graph/include/strided_slice_inst.h @@ -6,7 +6,6 @@ #include "intel_gpu/primitives/strided_slice.hpp" #include "primitive_inst.h" -#include "intel_gpu/runtime/error_handler.hpp" #include #include diff --git a/src/plugins/intel_gpu/src/graph/include/to_string_utils.h b/src/plugins/intel_gpu/src/graph/include/to_string_utils.h index 76cb4d3e69e..76cf9bc686f 100644 --- a/src/plugins/intel_gpu/src/graph/include/to_string_utils.h +++ b/src/plugins/intel_gpu/src/graph/include/to_string_utils.h @@ -21,12 +21,6 @@ namespace cldnn { -inline std::string to_string_hex(int val) { - std::stringstream stream; - stream << "0x" << std::uppercase << std::hex << val; - return stream.str(); -} - inline std::string bool_to_str(bool cond) { return cond ? "true" : "false"; } inline std::string get_extr_type(const std::string& str) { @@ -49,16 +43,6 @@ inline std::string fmt_to_str(format fmt) { inline std::string type_to_str(std::shared_ptr primitive) { return primitive->type_string(); } -inline std::string allocation_type_to_str(allocation_type type) { - switch (type) { - case allocation_type::cl_mem: return "cl_mem"; - case allocation_type::usm_host: return "usm_host"; - case allocation_type::usm_shared: return "usm_shared"; - case allocation_type::usm_device: return "usm_device"; - default: return "unknown"; - } -} - inline std::string activation_type_to_str(activation_func activation) { switch (activation) { case activation_func::none: return "none"; diff --git a/src/plugins/intel_gpu/src/graph/input_layout.cpp b/src/plugins/intel_gpu/src/graph/input_layout.cpp index 1b7a0a0eb5b..d12a6d9a74c 100644 --- a/src/plugins/intel_gpu/src/graph/input_layout.cpp +++ b/src/plugins/intel_gpu/src/graph/input_layout.cpp @@ -11,7 +11,7 @@ #include namespace { -bool has_optimized_users(input_layout_node const& node) { +bool has_optimized_users(cldnn::input_layout_node const& node) { for (auto& user : node.get_users()) { if (user->can_be_optimized()) { return true; diff --git a/src/plugins/intel_gpu/src/graph/loop.cpp b/src/plugins/intel_gpu/src/graph/loop.cpp index f7369447290..17b853cbb1d 100644 --- a/src/plugins/intel_gpu/src/graph/loop.cpp +++ b/src/plugins/intel_gpu/src/graph/loop.cpp @@ -8,6 +8,7 @@ #include "intel_gpu/primitives/data.hpp" #include "intel_gpu/primitives/mutable_data.hpp" #include "intel_gpu/graph/serialization/loop_serializer.hpp" +#include "intel_gpu/runtime/error_handler.hpp" #include #include #include diff --git a/src/plugins/intel_gpu/src/graph/lstm_gemm.cpp b/src/plugins/intel_gpu/src/graph/lstm_gemm.cpp index 992930841b3..a26d15d7cec 100644 --- a/src/plugins/intel_gpu/src/graph/lstm_gemm.cpp +++ b/src/plugins/intel_gpu/src/graph/lstm_gemm.cpp @@ -1,9 +1,9 @@ // Copyright (C) 2018-2023 Intel Corporation // SPDX-License-Identifier: Apache-2.0 // +#include "intel_gpu/runtime/error_handler.hpp" #include "lstm_gemm_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/matrix_nms.cpp b/src/plugins/intel_gpu/src/graph/matrix_nms.cpp index 24f58da7f3f..dd99b9230d7 100644 --- a/src/plugins/intel_gpu/src/graph/matrix_nms.cpp +++ b/src/plugins/intel_gpu/src/graph/matrix_nms.cpp @@ -65,6 +65,8 @@ std::string matrix_nms_inst::to_string(const matrix_nms_node& node) { } // namespace cldnn namespace ov { +using cldnn::matrix_nms; + template <> EnumNames& EnumNames::get() { static auto enum_names = EnumNames( diff --git a/src/plugins/intel_gpu/src/graph/multiclass_nms.cpp b/src/plugins/intel_gpu/src/graph/multiclass_nms.cpp index 394ea4bac47..8d451dce54d 100644 --- a/src/plugins/intel_gpu/src/graph/multiclass_nms.cpp +++ b/src/plugins/intel_gpu/src/graph/multiclass_nms.cpp @@ -4,7 +4,6 @@ #include -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include "multiclass_nms_inst.h" #include "primitive_type_base.h" diff --git a/src/plugins/intel_gpu/src/graph/mutable_data.cpp b/src/plugins/intel_gpu/src/graph/mutable_data.cpp index 6dc0af9d4bc..62bb16c3162 100644 --- a/src/plugins/intel_gpu/src/graph/mutable_data.cpp +++ b/src/plugins/intel_gpu/src/graph/mutable_data.cpp @@ -4,8 +4,6 @@ #include "mutable_data_inst.h" #include "primitive_type_base.h" #include "intel_gpu/runtime/memory.hpp" -#include -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include #include diff --git a/src/plugins/intel_gpu/src/graph/network.cpp b/src/plugins/intel_gpu/src/graph/network.cpp index bf30cb3bc7a..40217538874 100644 --- a/src/plugins/intel_gpu/src/graph/network.cpp +++ b/src/plugins/intel_gpu/src/graph/network.cpp @@ -13,26 +13,28 @@ #include "intel_gpu/runtime/stream.hpp" #include "intel_gpu/runtime/debug_configuration.hpp" #include "intel_gpu/runtime/half.hpp" +#include "intel_gpu/runtime/itt.hpp" #include "intel_gpu/graph/program.hpp" #include "intel_gpu/graph/network.hpp" #include "intel_gpu/graph/serialization/map_serializer.hpp" -#include "assign_inst.h" -#include "read_value_inst.h" -#include "reshape_inst.h" -#include "to_string_utils.h" #include "primitive_inst.h" #include "input_layout_inst.h" #include "mutable_data_inst.h" #include "condition_inst.h" #include "loop_inst.h" -#include "kernel_selector_helper.h" +#include "assign_inst.h" +#include "read_value_inst.h" +#include "reshape_inst.h" #include "program_helpers.h" -#include "intel_gpu/runtime/itt.hpp" +#include "to_string_utils.h" #include "kernels_cache.hpp" #include "compilation_context.hpp" +// TODO: Remove once we have an abstraction for kernels_cache +#include "kernel_base.h" + #include #include #include diff --git a/src/plugins/intel_gpu/src/graph/nodes_ordering.cpp b/src/plugins/intel_gpu/src/graph/nodes_ordering.cpp index d2d53c9c0ee..8c0dbb9c389 100644 --- a/src/plugins/intel_gpu/src/graph/nodes_ordering.cpp +++ b/src/plugins/intel_gpu/src/graph/nodes_ordering.cpp @@ -4,7 +4,6 @@ #include "intel_gpu/graph/program.hpp" #include "program_node.h" -#include "intel_gpu/runtime/error_handler.hpp" #include #include #include diff --git a/src/plugins/intel_gpu/src/graph/non_zero.cpp b/src/plugins/intel_gpu/src/graph/non_zero.cpp index 393e848f75f..01a9688ca9b 100644 --- a/src/plugins/intel_gpu/src/graph/non_zero.cpp +++ b/src/plugins/intel_gpu/src/graph/non_zero.cpp @@ -1,10 +1,10 @@ // Copyright (C) 2022 Intel Corporation // SPDX-License-Identifier: Apache-2.0 // + #include "non_zero_inst.h" #include "primitive_type_base.h" #include "intel_gpu/runtime/memory.hpp" -#include "intel_gpu/runtime/error_handler.hpp" #include "intel_gpu/runtime/debug_configuration.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/pass_manager.cpp b/src/plugins/intel_gpu/src/graph/pass_manager.cpp index 23dccf45034..22c8b9f5b44 100644 --- a/src/plugins/intel_gpu/src/graph/pass_manager.cpp +++ b/src/plugins/intel_gpu/src/graph/pass_manager.cpp @@ -13,6 +13,8 @@ #include #include +using namespace cldnn; + pass_manager::pass_manager(program& p) { pass_count = 0; auto path = get_dir_path(p.get_config()); diff --git a/src/plugins/intel_gpu/src/graph/primitive_inst.cpp b/src/plugins/intel_gpu/src/graph/primitive_inst.cpp index 2eb6b3ec55a..a6f7755d908 100644 --- a/src/plugins/intel_gpu/src/graph/primitive_inst.cpp +++ b/src/plugins/intel_gpu/src/graph/primitive_inst.cpp @@ -32,6 +32,7 @@ #include #include +namespace cldnn { namespace { bool is_optimized_output_user(const program_node* user) { @@ -80,7 +81,6 @@ bool is_user_cpu(const program_node* user) { } } // namespace -namespace cldnn { bool is_any_user_cpu(const std::list& users) { for (const auto& user : users) { if (is_user_cpu(user)) diff --git a/src/plugins/intel_gpu/src/graph/program.cpp b/src/plugins/intel_gpu/src/graph/program.cpp index db9c4436d0f..a455b061868 100644 --- a/src/plugins/intel_gpu/src/graph/program.cpp +++ b/src/plugins/intel_gpu/src/graph/program.cpp @@ -2,15 +2,14 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "intel_gpu/runtime/error_handler.hpp" #include "intel_gpu/runtime/memory.hpp" #include "intel_gpu/runtime/engine.hpp" #include "intel_gpu/runtime/debug_configuration.hpp" +#include "intel_gpu/runtime/itt.hpp" #include "intel_gpu/graph/program.hpp" #include -#include "kernel_selector_helper.h" #include "auto_tuner.h" #include "layout_optimizer.h" #include "pass_manager.h" @@ -66,8 +65,11 @@ #include "loop_inst.h" #include "reverse_inst.h" #include "to_string_utils.h" -#include "intel_gpu/runtime/itt.hpp" + +// TODO: Remove once we have interface for kernels cache #include "runtime/kernels_cache.hpp" + +// TODO: implement self-registration for impls #include "impls/ocl/register.hpp" #include "impls/cpu/register.hpp" #include "impls/common/register.hpp" @@ -96,6 +98,7 @@ #include #endif +using namespace cldnn; using namespace ov::intel_gpu; program::program(engine& engine_ref, diff --git a/src/plugins/intel_gpu/src/graph/program_dump_graph.cpp b/src/plugins/intel_gpu/src/graph/program_dump_graph.cpp index 7602a351ec7..a8cdbd1ac63 100644 --- a/src/plugins/intel_gpu/src/graph/program_dump_graph.cpp +++ b/src/plugins/intel_gpu/src/graph/program_dump_graph.cpp @@ -160,7 +160,7 @@ void dump_graph_init(std::ofstream& graph, return invalid_layout_msg; auto output_layout = ptr->get_output_layout(); - std::string out = fmt_to_str(output_layout.format); + std::string out = output_layout.format.to_string(); return out; }; diff --git a/src/plugins/intel_gpu/src/graph/program_node.cpp b/src/plugins/intel_gpu/src/graph/program_node.cpp index ce552b79a57..397b6bd3b36 100644 --- a/src/plugins/intel_gpu/src/graph/program_node.cpp +++ b/src/plugins/intel_gpu/src/graph/program_node.cpp @@ -126,7 +126,7 @@ std::unique_ptr program_node::desc_to_json() const { fused_node_info.add("dep start_idx", fused_desc.dep_start_idx); json_composite info; info.add("data type", dt_to_str(fused_desc.output_layout.data_type)); - info.add("format", fmt_to_str(output_layouts[0].format)); + info.add("format", output_layouts[0].format.to_string()); info.add("size", output_layouts[0].to_short_string()); fused_node_info.add("output layout", info); fused_nodes_info.add("fused primitive idx " + std::to_string(index++), fused_node_info); diff --git a/src/plugins/intel_gpu/src/graph/pyramid_roi_align.cpp b/src/plugins/intel_gpu/src/graph/pyramid_roi_align.cpp index 0a61256c19f..c7979ad98a6 100644 --- a/src/plugins/intel_gpu/src/graph/pyramid_roi_align.cpp +++ b/src/plugins/intel_gpu/src/graph/pyramid_roi_align.cpp @@ -3,7 +3,6 @@ // #include "pyramid_roi_align_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/quantize.cpp b/src/plugins/intel_gpu/src/graph/quantize.cpp index 482ca268ad3..dad0b12bf96 100644 --- a/src/plugins/intel_gpu/src/graph/quantize.cpp +++ b/src/plugins/intel_gpu/src/graph/quantize.cpp @@ -6,7 +6,6 @@ #include "binary_convolution_inst.h" #include "primitive_type_base.h" #include "intel_gpu/runtime/memory.hpp" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include "data_inst.h" #include diff --git a/src/plugins/intel_gpu/src/graph/read_value.cpp b/src/plugins/intel_gpu/src/graph/read_value.cpp index 6d20c3ec103..af4023bfee2 100644 --- a/src/plugins/intel_gpu/src/graph/read_value.cpp +++ b/src/plugins/intel_gpu/src/graph/read_value.cpp @@ -6,7 +6,6 @@ #include "primitive_type_base.h" #include #include -#include namespace cldnn { GPU_DEFINE_PRIMITIVE_TYPE_ID(read_value) diff --git a/src/plugins/intel_gpu/src/graph/reduce.cpp b/src/plugins/intel_gpu/src/graph/reduce.cpp index 6ecf6763bf7..cb5d82c8758 100644 --- a/src/plugins/intel_gpu/src/graph/reduce.cpp +++ b/src/plugins/intel_gpu/src/graph/reduce.cpp @@ -5,9 +5,7 @@ #include "reduce_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" -#include "data_inst.h" #include #include diff --git a/src/plugins/intel_gpu/src/graph/resample.cpp b/src/plugins/intel_gpu/src/graph/resample.cpp index 5ecd304251b..a0db38603af 100644 --- a/src/plugins/intel_gpu/src/graph/resample.cpp +++ b/src/plugins/intel_gpu/src/graph/resample.cpp @@ -3,7 +3,6 @@ // #include "resample_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include #include "json_object.h" diff --git a/src/plugins/intel_gpu/src/graph/reverse_sequence.cpp b/src/plugins/intel_gpu/src/graph/reverse_sequence.cpp index 8f0bc73c1dd..715b98f81c0 100644 --- a/src/plugins/intel_gpu/src/graph/reverse_sequence.cpp +++ b/src/plugins/intel_gpu/src/graph/reverse_sequence.cpp @@ -5,7 +5,6 @@ #include "reverse_sequence_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/roi_align.cpp b/src/plugins/intel_gpu/src/graph/roi_align.cpp index d59bbe3e5ea..825ba7c3d87 100644 --- a/src/plugins/intel_gpu/src/graph/roi_align.cpp +++ b/src/plugins/intel_gpu/src/graph/roi_align.cpp @@ -46,6 +46,8 @@ std::string roi_align_inst::to_string(roi_align_node const& node) { } // namespace cldnn namespace ov { +using cldnn::roi_align; + template <> EnumNames& EnumNames::get() { static auto enum_names = EnumNames("PoolingMode", {{"max", roi_align::PoolingMode::max}, diff --git a/src/plugins/intel_gpu/src/graph/roi_pooling.cpp b/src/plugins/intel_gpu/src/graph/roi_pooling.cpp index 845e2f4a505..4e887374674 100644 --- a/src/plugins/intel_gpu/src/graph/roi_pooling.cpp +++ b/src/plugins/intel_gpu/src/graph/roi_pooling.cpp @@ -4,7 +4,6 @@ #include "roi_pooling_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/scatter_elements_update.cpp b/src/plugins/intel_gpu/src/graph/scatter_elements_update.cpp index 8c53e91d434..5c12acafd59 100644 --- a/src/plugins/intel_gpu/src/graph/scatter_elements_update.cpp +++ b/src/plugins/intel_gpu/src/graph/scatter_elements_update.cpp @@ -2,10 +2,10 @@ // SPDX-License-Identifier: Apache-2.0 // +#include "intel_gpu/runtime/error_handler.hpp" #include "scatter_elements_update_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/scatter_nd_update.cpp b/src/plugins/intel_gpu/src/graph/scatter_nd_update.cpp index 0ad55d88571..4d11b4fb8e5 100644 --- a/src/plugins/intel_gpu/src/graph/scatter_nd_update.cpp +++ b/src/plugins/intel_gpu/src/graph/scatter_nd_update.cpp @@ -5,7 +5,6 @@ #include "scatter_nd_update_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/scatter_update.cpp b/src/plugins/intel_gpu/src/graph/scatter_update.cpp index b1e6cd954ab..07c8e0ab8f9 100644 --- a/src/plugins/intel_gpu/src/graph/scatter_update.cpp +++ b/src/plugins/intel_gpu/src/graph/scatter_update.cpp @@ -5,7 +5,6 @@ #include "scatter_update_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/graph/shape_of.cpp b/src/plugins/intel_gpu/src/graph/shape_of.cpp index 88742eeaeb9..86081cd9c38 100644 --- a/src/plugins/intel_gpu/src/graph/shape_of.cpp +++ b/src/plugins/intel_gpu/src/graph/shape_of.cpp @@ -4,7 +4,6 @@ #include "shape_of_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" #include "to_string_utils.h" #include diff --git a/src/plugins/intel_gpu/src/graph/space_to_depth.cpp b/src/plugins/intel_gpu/src/graph/space_to_depth.cpp index 55f9c17b58c..137a59d68a8 100644 --- a/src/plugins/intel_gpu/src/graph/space_to_depth.cpp +++ b/src/plugins/intel_gpu/src/graph/space_to_depth.cpp @@ -79,7 +79,7 @@ std::string space_to_depth_inst::to_string(space_to_depth_node const& node) { std::stringstream primitive_description; - std::string depth_mode = (desc->mode == static_cast (kernel_selector::SpaceToDepthMode::BLOCKS_FIRST)) ? + std::string depth_mode = (desc->mode == cldnn::space_to_depth::depth_mode::blocks_first) ? "blocks_first" : "depth_first"; diff --git a/src/plugins/intel_gpu/src/graph/strided_slice.cpp b/src/plugins/intel_gpu/src/graph/strided_slice.cpp index e94d1790139..7d13b4e5b42 100644 --- a/src/plugins/intel_gpu/src/graph/strided_slice.cpp +++ b/src/plugins/intel_gpu/src/graph/strided_slice.cpp @@ -4,9 +4,7 @@ #include "strided_slice_inst.h" #include "primitive_type_base.h" -#include "intel_gpu/runtime/error_handler.hpp" #include "json_object.h" -#include "data_inst.h" #include #include diff --git a/src/plugins/intel_gpu/src/graph/tile.cpp b/src/plugins/intel_gpu/src/graph/tile.cpp index d7ca41dfbba..ba961c744d4 100644 --- a/src/plugins/intel_gpu/src/graph/tile.cpp +++ b/src/plugins/intel_gpu/src/graph/tile.cpp @@ -7,7 +7,6 @@ #include "primitive_type_base.h" #include "intel_gpu/runtime/memory.hpp" -#include "intel_gpu/runtime/error_handler.hpp" #include "intel_gpu/runtime/format.hpp" #include "json_object.h" #include diff --git a/src/plugins/intel_gpu/src/plugin/program.cpp b/src/plugins/intel_gpu/src/plugin/program.cpp index 55bd9102883..8edaf998f76 100644 --- a/src/plugins/intel_gpu/src/plugin/program.cpp +++ b/src/plugins/intel_gpu/src/plugin/program.cpp @@ -380,7 +380,7 @@ std::shared_ptr Program::BuildProgram(const std::vector diff --git a/src/plugins/intel_gpu/tests/test_utils/network_test.h b/src/plugins/intel_gpu/tests/test_utils/network_test.h index 7aaed5dd076..78ef454fc00 100644 --- a/src/plugins/intel_gpu/tests/test_utils/network_test.h +++ b/src/plugins/intel_gpu/tests/test_utils/network_test.h @@ -22,6 +22,7 @@ #include #include +using namespace cldnn; namespace tests { diff --git a/src/plugins/intel_gpu/tests/test_utils/test_utils.h b/src/plugins/intel_gpu/tests/test_utils/test_utils.h index d46935aee53..b8c8cc5addd 100644 --- a/src/plugins/intel_gpu/tests/test_utils/test_utils.h +++ b/src/plugins/intel_gpu/tests/test_utils/test_utils.h @@ -68,7 +68,7 @@ bool has_node_with_type(cldnn::program& prog) { return false; } -inline bool has_node(cldnn::program& prog, primitive_id id) { +inline bool has_node(cldnn::program& prog, cldnn::primitive_id id) { for (auto node : prog.get_processing_order()) { if (node->id() == id) return true; @@ -409,7 +409,7 @@ public: test_params() : fmt(cldnn::format::bfyx) { } - test_params(cldnn::data_types dt, cldnn::format input_format, int32_t batch_size, int32_t feature_size, cldnn::tensor input_size, ExecutionConfig config = {}) : + test_params(cldnn::data_types dt, cldnn::format input_format, int32_t batch_size, int32_t feature_size, cldnn::tensor input_size, cldnn::ExecutionConfig config = {}) : data_type(dt), fmt(input_format), network_config(config) { @@ -423,7 +423,7 @@ public: void * opaque_custom_param = nullptr; - ExecutionConfig network_config; + cldnn::ExecutionConfig network_config; std::string print(); static std::string print_tensor(cldnn::tensor tensor); @@ -574,42 +574,42 @@ T div_up(const T a, const U b) { } template -std::vector get_output_values_to_float(network& net, const primitive_id& output_id, size_t max_cnt = std::numeric_limits::max()) { +std::vector get_output_values_to_float(cldnn::network& net, const cldnn::primitive_id& output_id, size_t max_cnt = std::numeric_limits::max()) { std::vector ret; auto ptr = net.get_output_memory(output_id); auto out_ids = net.get_output_ids(); if (find(out_ids.begin(), out_ids.end(), output_id) == out_ids.end()) IE_THROW() << "Non output node's memory may have been reused. " "Make target node to output by using ov::intel_gpu::custom_outputs in ExecutionConfig."; - mem_lock mem(ptr, net.get_stream()); - if (ptr->get_layout().data_type != type_to_data_type::value) - IE_THROW() << "target type " << data_type_traits::name(type_to_data_type::value) - << " mismatched with actual type " << data_type_traits::name(ptr->get_layout().data_type); + cldnn::mem_lock mem(ptr, net.get_stream()); + if (ptr->get_layout().data_type != cldnn::type_to_data_type::value) + IE_THROW() << "target type " << cldnn::data_type_traits::name(cldnn::type_to_data_type::value) + << " mismatched with actual type " << cldnn::data_type_traits::name(ptr->get_layout().data_type); for (size_t i = 0; i < std::min(max_cnt, ptr->get_layout().count()); i++) ret.push_back(mem[i]); return ret; } -inline std::vector get_output_values_to_float(network& net, const primitive_id& output_id, size_t max_cnt = std::numeric_limits::max()) { +inline std::vector get_output_values_to_float(cldnn::network& net, const cldnn::primitive_id& output_id, size_t max_cnt = std::numeric_limits::max()) { switch(net.get_output_layout(output_id).data_type){ - case data_types::f16: + case cldnn::data_types::f16: return get_output_values_to_float(net, output_id, max_cnt); - case data_types::f32: + case cldnn::data_types::f32: return get_output_values_to_float(net, output_id, max_cnt); - case data_types::i8: + case cldnn::data_types::i8: return get_output_values_to_float(net, output_id, max_cnt); - case data_types::u8: + case cldnn::data_types::u8: return get_output_values_to_float(net, output_id, max_cnt); - case data_types::i32: + case cldnn::data_types::i32: return get_output_values_to_float(net, output_id, max_cnt); - case data_types::i64: + case cldnn::data_types::i64: return get_output_values_to_float(net, output_id, max_cnt); default: IE_THROW() << "Unknown output data_type"; } } -double default_tolerance(data_types dt); +double default_tolerance(cldnn::data_types dt); // inline void print_bin_blob(cldnn::memory& mem, std::string name) // { // auto&& size = mem.get_layout().get_tensor(); @@ -744,16 +744,16 @@ inline cldnn::network::ptr get_network(cldnn::engine& engine, cldnn::network::ptr network; if (is_caching_test) { std::cout << "cached" << std::endl; - membuf mem_buf; + cldnn::membuf mem_buf; { cldnn::network _network(engine, topology, config); std::ostream out_mem(&mem_buf); - BinaryOutputBuffer ob = BinaryOutputBuffer(out_mem); + cldnn::BinaryOutputBuffer ob = cldnn::BinaryOutputBuffer(out_mem); _network.save(ob); } { std::istream in_mem(&mem_buf); - BinaryInputBuffer ib = BinaryInputBuffer(in_mem, engine); + cldnn::BinaryInputBuffer ib = cldnn::BinaryInputBuffer(in_mem, engine); network = std::make_shared(ib, config, stream, engine); } } else { diff --git a/src/plugins/intel_gpu/tests/transformations/decompose_reduce_for_false_keepdims_test.cpp b/src/plugins/intel_gpu/tests/transformations/decompose_reduce_for_false_keepdims_test.cpp index ddfaa9f89f9..5d470e1e8d0 100644 --- a/src/plugins/intel_gpu/tests/transformations/decompose_reduce_for_false_keepdims_test.cpp +++ b/src/plugins/intel_gpu/tests/transformations/decompose_reduce_for_false_keepdims_test.cpp @@ -21,6 +21,7 @@ using namespace testing; using namespace ::tests; +using namespace cldnn; using InputShape = ngraph::PartialShape; using KeepDims = bool;