[GPU] Code cleanup (#16014)

* [GPU] Improve exception message for program build

* [GPU] Code cleanup

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 <string>
 

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 <string>

src/plugins/intel_gpu/src/graph/assign.cpp

@@ -2,11 +2,10 @@
 // SPDX-License-Identifier: Apache-2.0
 //
 
-#include <assign_inst.h>
+#include "assign_inst.h"
 #include "primitive_type_base.h"
 #include <sstream>
 #include <json_object.h>
-#include <data_inst.h>
 
 namespace cldnn {
 GPU_DEFINE_PRIMITIVE_TYPE_ID(assign)

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 <string>
 #include <vector>
 

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"

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 <string>
 

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)

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 <string>
 

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 <string>
 

src/plugins/intel_gpu/src/graph/experimental_detectron_detection_output.cpp

@@ -5,7 +5,6 @@
 #include <string>
 
 #include "experimental_detectron_detection_output_inst.hpp"
-#include "intel_gpu/runtime/error_handler.hpp"
 #include "json_object.h"
 #include "primitive_type_base.h"
 

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 <string>
 

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 <string>
 

src/plugins/intel_gpu/src/graph/experimental_detectron_topk_rois.cpp

@@ -6,7 +6,6 @@
 #include "primitive_type_base.h"
 #include <sstream>
 #include <json_object.h>
-#include <data_inst.h>
 
 namespace cldnn {
 GPU_DEFINE_PRIMITIVE_TYPE_ID(experimental_detectron_topk_rois)

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 <string>
 

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 <string>
 #include <algorithm>

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 <string>
 

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 <string>
 

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 <string>
 

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 <string>
 #include <utility>

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 <string>
 

src/plugins/intel_gpu/src/graph/graph_optimizer/add_required_reorders.cpp

@@ -13,6 +13,8 @@
 #include <vector>
 #include <stdexcept>
 
+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.

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 <iomanip>
 #include <string>

src/plugins/intel_gpu/src/graph/graph_optimizer/post_input_reorder.cpp

@@ -8,6 +8,8 @@
 #include <memory>
 #include <stdexcept>
 
+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.

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 <vector>
 #include <list>
 #include <memory>
 #include <string>
 #include <utility>
-#include "intel_gpu/runtime/error_handler.hpp"
+
+using namespace cldnn;
 
 void pre_replace_deconv::run(program& p) {
     bool update_processing_order = false;

src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_primitive_fusing.cpp

@@ -47,11 +47,12 @@
 #include <string>
 #include <utility>
 #include <deque>
-#include "intel_gpu/runtime/error_handler.hpp"
 #ifdef ENABLE_ONEDNN_FOR_GPU
 #include <impls/onednn/utils.hpp>
 #endif
 
+using namespace cldnn;
+
 void prepare_primitive_fusing::run(program& p) {
     fuse_reorders(p);
     remove_redundant_reshape(p);

src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_quantization.cpp

@@ -19,6 +19,8 @@
 #include <memory>
 #include <vector>
 
+using namespace cldnn;
+
 template<typename T>
 bool check_binarization(memory::ptr mem_input_low, memory::ptr mem_input_high, program& p) {
     bool is_binarization = true;

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 <algorithm>

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 <vector>
 #include <algorithm>
 

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 <vector>
@@ -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() {}
 

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 {

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 <vector>
 #include <utility>
 
-#ifdef FIX_OPENMP_RELEASE_ISSUE
-#ifdef OPENMP_FOUND
-#include <omp.h>
-#endif
-#endif
-
 namespace cldnn {
 namespace cpu {
 
-namespace {
-    using bounding_box = cldnn::cpu::bounding_box;
-}  // namespace
-
 template <typename T>
 bool comp_score_descend(const std::pair<float, T>& pair1,
                         const std::pair<float, T>& pair2) {
@@ -296,15 +286,6 @@ public:
             std::vector<std::vector<std::pair<float, int>>>& conf_per_image = confidences[image];
             std::map<int, std::vector<int>> 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<int>(args->num_classes); ++cls) {
                     if (static_cast<int>(cls) == args->background_label_id) {

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 <vector>
 #include <queue>

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"
 

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 {

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<kernel_selector::base_activation_params>& params) {
+inline void convert_new_activation_func(const cldnn::activation& prim, std::vector<kernel_selector::base_activation_params>& params) {
     params.insert(params.begin(), {get_kernel_selector_activation_param(prim.activation_function),
                                    prim.additional_params.a,
                                    prim.additional_params.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 {

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"
 

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 <vector>
-
-using namespace cldnn;
 
 namespace cldnn {
 namespace ocl {

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 <algorithm>
-#include <memory>
 
 namespace cldnn {
 namespace ocl {

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 {

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 {

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 {

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 <initializer_list>
-
 namespace cldnn {
 namespace ocl {
 

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 <vector>
-
-using namespace cldnn;
 
 namespace cldnn {
 namespace ocl {

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 <algorithm>
-#include <memory>
 
 namespace cldnn {
 namespace ocl {

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 {

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 <algorithm>
-
-using namespace cldnn;
-
 namespace cldnn {
 namespace ocl {
 

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 {

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 {
 

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 <map>
 #include <sstream>
@@ -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)));
 

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 <algorithm>
 
 namespace cldnn {
 namespace ocl {

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 <algorithm>
 
 namespace cldnn {
 namespace ocl {

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 {

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 <vector>
 
 namespace cldnn {
 namespace ocl {

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

@@ -2,16 +2,12 @@
 // SPDX-License-Identifier: Apache-2.0
 //
 
-#include <dft/dft_kernel_ref.h>
-#include <dft/dft_kernel_selector.h>
-#include <dft_inst.h>
-#include <kernel_selector_helper.h>
-
-#include <impls/implementation_map.hpp>
-#include <intel_gpu/runtime/error_handler.hpp>
-
 #include "primitive_base.hpp"
 
+#include "dft_inst.h"
+#include "dft/dft_kernel_ref.h"
+#include "dft/dft_kernel_selector.h"
+
 namespace cldnn {
 namespace ocl {
 

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 <vector>
 
 namespace cldnn {
 namespace ocl {

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 {

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 {

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

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 <ed_pgg/prior_grid_generator_kernel_ref.h>
-#include <ed_pgg/prior_grid_generator_kernel_selector.h>
-#include <experimental_detectron_prior_grid_generator_inst.h>
-#include <kernel_selector_helper.h>
-
-#include <impls/implementation_map.hpp>
-#include <intel_gpu/runtime/error_handler.hpp>
-
 #include "primitive_base.hpp"
 
+#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 {
 

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"
 

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

@@ -2,13 +2,11 @@
 // SPDX-License-Identifier: Apache-2.0
 //
 
-#include <experimental_detectron_topk_rois_inst.h>
-#include "intel_gpu/runtime/error_handler.hpp"
-#include <impls/implementation_map.hpp>
-#include <ed_topkroi/topk_rois_ref.h>
-#include <ed_topkroi/topk_rois_kernel_selector.h>
 #include "primitive_base.hpp"
-#include <vector>
+
+#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 {

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"
 

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

@@ -2,19 +2,12 @@
 // SPDX-License-Identifier: Apache-2.0
 //
 
-#include <data_inst.h>
-#include <eye/eye_kernel_ref.h>
-#include <eye/eye_kernel_selector.h>
-#include <eye_inst.h>
-
-#include <algorithm>
-#include <cstddef>
-#include <impls/implementation_map.hpp>
-#include <intel_gpu/runtime/error_handler.hpp>
-#include <vector>
-
 #include "primitive_base.hpp"
 
+#include "eye_inst.h"
+#include "eye/eye_kernel_ref.h"
+#include "eye/eye_kernel_selector.h"
+
 namespace cldnn {
 namespace ocl {
 

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 <memory>
-#include <algorithm>
-
 namespace cldnn {
 namespace ocl {
 

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 {

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 {

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 {
 

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 {

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 <algorithm>
-#include "kernel_selector_helper.h"
-#include "primitive_base.hpp"
 
 namespace cldnn {
 namespace ocl {

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

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 <vector>
 
 namespace cldnn {
 namespace ocl {

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 {

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 <algorithm>
-
-using namespace cldnn;
-
 namespace cldnn {
 namespace ocl {
 

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

@@ -34,7 +34,10 @@
 #include <string>
 #include <vector>
 
+
 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

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 <vector>
 #include <memory>
 
-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<const primitive> desc;
-    size_t unique_id;
-    std::vector<layout> input_layouts;
-    std::vector<layout> output_layouts;
-    std::vector<tensor> input_offsets;
-    std::vector<cldnn::fused_primitive_desc> fused_desc;
-#ifdef ENABLE_ONEDNN_FOR_GPU
-    std::vector<cldnn::fused_primitive_desc_onednn> 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<size_t, memory::ptr> memory_deps = {};
-    size_t primary_input_idx = 0;
-
-    memory::ptr reordered_weights = nullptr;
-
-    kernel_impl_params() {}
-
-    kernel_impl_params(program& _prog,
-                       std::shared_ptr<const primitive> _desc,
-                       size_t _uid,
-                       const std::vector<layout>& _in_layouts,
-                       const std::vector<layout>& _out_layouts,
-                       const std::vector<cldnn::fused_primitive_desc>& _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 <class PType>
-    std::shared_ptr<const PType> typed_desc() const { return std::static_pointer_cast<const PType>(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 <typename T = std::uint32_t>
 kernel_selector::dim_tensor<T> 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 <typename optional_params_t>
 inline optional_params_t get_default_weights_bias_optional_params(const program& program) {
     return get_default_optional_params<optional_params_t>(program);
 }
+}  // namespace cldnn

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"
 

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 {

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 {

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<primitive_impl> 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<kernel_selector::lstm_elt_params>(impl_param);
-        auto lstm_elt_optional_params =
-            get_default_optional_params<kernel_selector::lstm_elt_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<lstm_elt>();
+        auto params = get_default_params<kernel_selector::lstm_elt_params>(impl_param);
+        auto optional_params = get_default_optional_params<kernel_selector::lstm_elt_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<int32_t>(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<int32_t>(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<lstm_elt_impl>(best_kernel);
+        return {params, optional_params};
     }
 };
 
 namespace detail {
 
 attach_lstm_elt_impl::attach_lstm_elt_impl() {
-    implementation_map<lstm_elt>::add(impl_types::ocl, lstm_elt_impl::create, {
+    implementation_map<lstm_elt>::add(impl_types::ocl, typed_primitive_impl_ocl<lstm_elt>::create<lstm_elt_impl>, {
         std::make_tuple(data_types::f32, format::bfyx),
         std::make_tuple(data_types::f16, format::bfyx),
         std::make_tuple(data_types::f32, format::fyxb),

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 {

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 {

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 {

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 {

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 <algorithm>
-
-using namespace cldnn;
-
 namespace cldnn {
 namespace ocl {
 

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 {

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 {

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 <algorithm>
-
-using namespace cldnn;
-
 namespace cldnn {
 namespace ocl {
 

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 <vector>
 
 namespace cldnn {
 namespace ocl {

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 {
 

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 <algorithm>
 
 namespace cldnn {
 namespace ocl {

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

@@ -4,11 +4,6 @@
 
 #pragma once
 
-#include <thread>
-#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 <vector>
 #include <list>
 #include <utility>

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

@@ -2,16 +2,12 @@
 // SPDX-License-Identifier: Apache-2.0
 //
 
-#include <prior_box/prior_box_kernel_ref.h>
-#include <prior_box/prior_box_kernel_selector.h>
-#include <prior_box_inst.h>
-
-#include <impls/implementation_map.hpp>
-#include <vector>
-
-#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 {
 

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 <cmath>
 

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 {

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

@@ -2,13 +2,11 @@
 // SPDX-License-Identifier: Apache-2.0
 //
 
-#include <random_uniform_inst.h>
-#include <random_uniform/random_uniform_kernel_ref.h>
-#include "intel_gpu/runtime/error_handler.hpp"
-#include <impls/implementation_map.hpp>
-#include <random_uniform/random_uniform_kernel_selector.h>
 #include "primitive_base.hpp"
-#include <vector>
+
+#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 {

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

@@ -2,13 +2,11 @@
 // SPDX-License-Identifier: Apache-2.0
 //
 
-#include <range_inst.h>
 #include "primitive_base.hpp"
-#include <impls/implementation_map.hpp>
-#include <kernel_selector_helper.h>
-#include <range/range_kernel_selector.h>
-#include <range/range_kernel_ref.h>
-#include <intel_gpu/runtime/error_handler.hpp>
+
+#include "range_inst.h"
+#include "range/range_kernel_selector.h"
+#include "range/range_kernel_ref.h"
 
 namespace cldnn {
 namespace ocl {

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 {

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 {

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 {