[GPU] Add fusion through feature (#9674)

* [GPU] Add fuse through feature

* Apply review comments


Tasks marked as not passed:
* https://dev.azure.com/openvinoci/dldt/_build/results?buildId=313769&view=results
* https://dev.azure.com/openvinoci/dldt/_build/results?buildId=313766&view=results

src/plugins/intel_gpu/include/intel_gpu/graph/program.hpp

@@ -44,8 +44,9 @@ struct program {
     friend class prepare_conv_eltw_fusing;   // to be removed when possible
     friend class reorder_inputs;             // to be removed when possible
     friend class remove_redundant_reorders;  // to be removed when possible
-    friend class program_wrapper;       // this class is intended to extend the interface of program for
+    friend class program_wrapper;            // this class is intended to extend the interface of program for
                                              // the usage within tests_core_internal project only
+    friend class prepare_primitive_fusing_through;   // to be removed when possible
 public:
     struct nodes_ordering {
     public:
@@ -189,6 +190,12 @@ public:
     // returns if 'node' has been extracted and removed successfully
     bool extract_and_remove(program_node& node);
 
+    bool extract(program_node& node);
+
+    bool move_node(program_node& node,
+                   program_node& new_prev,
+                   program_node& new_next);
+
     // Fuses two nodes into fused_node and removes peer_node from graph
     void fuse_nodes(program_node& fused_node,
                     program_node& peer_node,
@@ -324,7 +331,7 @@ private:
 
     void rename(program_node& node, primitive_id const& new_id);
     void swap_names(program_node& node1, program_node& node2);
-    void replace_all_usages(program_node& old_node, program_node& new_node);
+    void replace_all_usages(program_node& old_node, program_node& new_node, bool remove_if_dangling = true);
 
     // old_node - node which will be replaced
     // new_node - node which will replace the old one

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

@@ -741,10 +741,7 @@ void prepare_primitive_fusing::fuse_simple_primitives(program &p) {
 
             should_fuse |= input_data.is_type<scale>();
 
-            // Here we need to check that Eltwise already has fused ops to avoid missing Activation primitive in
-            // case `Conv -> Eltwise -> Activation` which will be replaced via fused_conv_eltwise primitive later
-            // without handling any fused ops
-            should_fuse |= input_data.is_type<eltwise>() && eltwise_supports_fusings(input_data.as<eltwise>()) && input_data.has_fused_primitives();
+            should_fuse |= input_data.is_type<eltwise>() && eltwise_supports_fusings(input_data.as<eltwise>());
 
             if (!should_fuse)
                 return;
@@ -832,6 +829,14 @@ void prepare_primitive_fusing::fuse_simple_primitives(program &p) {
             auto out_dt_is_i8_u8 = data_type_traits::is_i8_u8(out_dt);
             auto in_dt_is_i8_u8 = data_type_traits::is_i8_u8(in_dt);
 
+            bool per_tensor_values = quantize_node.get_scale_shift_opt() &&
+                                     quantize_node.get_per_tensor_input_scale() &&
+                                     quantize_node.get_per_tensor_input_shift() &&
+                                     quantize_node.get_per_tensor_input_range() &&
+                                     quantize_node.get_per_tensor_output_scale() &&
+                                     quantize_node.get_per_tensor_output_shift() &&
+                                     quantize_node.get_per_tensor_output_range();
+
             bool should_fuse = input_data.is_type<binary_convolution>() &&
                                ((out_dt == data_types::bin &&
                                quantize_node.get_dependencies().size() == 5 &&
@@ -911,6 +916,11 @@ void prepare_primitive_fusing::fuse_simple_primitives(program &p) {
 
             should_fuse |= input_data.is_type<scale>() && quantize_node.get_scale_shift_opt();
 
+            should_fuse |= input_data.is_type<softmax>() &&
+                           input_data.as<softmax>().get_primitive()->dimension == softmax::dimension_t::normalize_f &&
+                           per_tensor_values;
+
+
             if (!should_fuse)
                 return;
 

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

@@ -0,0 +1,174 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+
+#include "intel_gpu/runtime/error_handler.hpp"
+#include "pass_manager.h"
+#include "program_helpers.h"
+#include "strided_slice_inst.h"
+#include "reshape_inst.h"
+#include "data_inst.h"
+#include "eltwise_inst.h"
+#include "mutable_data_inst.h"
+#include <vector>
+#include <memory>
+
+using namespace cldnn;
+
+void prepare_primitive_fusing_through::run(program& p) {
+    auto try_fuse_through = [&](program_node& node) -> std::vector<program_node*> {
+        // This function tries to fuse peer_node to first non reorder or reshape previous primitive.
+        // It returns chain of primitives (reshapes and reorders) including potential fused_node (e.g. Conv, FC, etc)
+        // at the end of vector.
+        // There are some limitations:
+        // 1. As for now we can fuse only through chain of reorders and reshapes primitives
+        // 2. There are three types of supported primitives:
+        //    - Quantize with per-tensor parameters
+        //    - Eltwise with single scale value in constant buffer
+        //    - Activation w/o additional buffers
+        // 3. Paddings are not allowed for fused node
+        auto can_raise_up_through = [](program_node* node) {
+            if (!node->is_type<reshape>() && !node->is_type<reorder>())
+                return false;
+
+            if (node->get_dependencies().empty())
+                return false;
+
+            if (node->get_users().size() > 1)
+                return false;
+
+            if (node->is_type<reorder>() &&
+                node->get_output_layout().data_type != node->get_dependency(0).get_output_layout().data_type)
+                return false;
+
+            return true;
+        };
+
+        std::vector<program_node*> pass_through;
+        program_node* fuse_through = &node;
+        pass_through.push_back(fuse_through);
+
+        bool can_raise_up = can_raise_up_through(fuse_through);
+        while (can_raise_up) {
+            fuse_through = &fuse_through->get_dependency(0);
+            can_raise_up = can_raise_up_through(fuse_through);
+            pass_through.push_back(fuse_through);
+        }
+
+        return pass_through;
+    };
+
+    auto node_itr = p.get_processing_order().begin();
+    while (node_itr != p.get_processing_order().end()) {
+        auto node = (*node_itr++);
+        cldnn::program_node* input_node;
+
+        if (node->is_output() || node->is_constant())
+            continue;
+
+        if (node->is_type<activation>()) {
+            if (node->get_dependencies().size() > 1)
+                continue;
+
+            input_node = &node->get_dependency(0);
+        } else if (node->is_type<quantize>()) {
+            auto& quantize_node = node->as<quantize>();
+            bool per_tensor_values = quantize_node.get_scale_shift_opt() &&
+                                     quantize_node.get_per_tensor_input_scale() &&
+                                     quantize_node.get_per_tensor_input_shift() &&
+                                     quantize_node.get_per_tensor_input_range() &&
+                                     quantize_node.get_per_tensor_output_scale() &&
+                                     quantize_node.get_per_tensor_output_shift() &&
+                                     quantize_node.get_per_tensor_output_range();
+
+            if (!per_tensor_values)
+                continue;
+
+            input_node = &node->get_dependency(0);
+        } else if (node->is_type<eltwise>()) {
+            if (node->get_dependencies().size() !=2)
+                continue;
+
+            size_t second_input_idx = 0;
+            bool has_constant_input = false;
+            for (size_t i = 0; i < node->get_dependencies().size(); i++) {
+                auto& dep = node->get_dependency(i);
+                if (dep.is_constant() && dep.get_output_layout().size == cldnn::tensor(1)) {
+                    second_input_idx = i ^ 1;
+                    has_constant_input = true;
+                    break;
+                }
+            }
+
+            if (!has_constant_input)
+                continue;
+
+            input_node = &node->get_dependency(second_input_idx);
+        } else {
+            continue;
+        }
+
+        auto fuse_through_order = try_fuse_through(*input_node);
+        bool use_fuse_through = fuse_through_order.size() > 1;
+        bool fused_node_has_multiple_users = fuse_through_order.back()->get_users().size() > 1;
+
+        if (!use_fuse_through || fused_node_has_multiple_users)
+            continue;
+
+        if (static_cast<bool>(node->get_output_layout().data_padding))
+            continue;
+
+        auto new_prev = fuse_through_order[fuse_through_order.size() - 1];
+        auto new_next = fuse_through_order[fuse_through_order.size() - 2];
+
+        if (new_prev->is_type<input_layout>() ||
+            new_prev->is_type<mutable_data>() ||
+            new_prev->is_type<quantize>())
+            continue;
+
+        std::vector<cldnn::program_node*> dependencies;
+        for (auto dep : node->get_dependencies()) {
+            if (dep == input_node)
+                continue;
+            dependencies.push_back(dep);
+        }
+
+        for (auto dep : dependencies)
+            p.remove_connection(*dep, *node);
+
+        p.move_node(*node, *new_prev, *new_next);
+
+        for (auto dep : dependencies)
+            p.add_connection(*dep, *node);
+
+        // Update node's layout and keep original data type
+        auto target_dt = node->get_output_layout().data_type;
+        layout new_node_layout = new_prev->get_output_layout();
+        new_node_layout.data_type = target_dt;
+        node->set_output_layout(new_node_layout, false);
+
+        // Update intermediate nodes
+        auto node_itr = std::next(fuse_through_order.rbegin());
+        while (node_itr != fuse_through_order.rend()) {
+            auto itermediate_node = *node_itr++;
+
+            if (itermediate_node->is_type<reorder>()) {
+                // We can't modify reorder's output type directly, so replace old node with new one
+                auto reorder_layout = itermediate_node->get_output_layout();
+                reorder_layout.data_type = target_dt;
+                auto r_prim = std::make_shared<reorder>(itermediate_node->id() + "_reorder_to_req_dt", itermediate_node->id(), reorder_layout);
+                auto& new_reorder = p.get_or_create(r_prim);
+
+                p.add_intermediate(new_reorder, *itermediate_node->get_users().front(), *itermediate_node);
+                new_reorder.recalc_output_layout(false);
+
+                p.add_optimized_primitive_info(itermediate_node->id());
+                p.extract_and_remove(*itermediate_node);
+            } else {
+                itermediate_node->recalc_output_layout(false);
+            }
+        }
+    }
+}

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

@@ -185,6 +185,12 @@ private:
     void conv_eltwise_read_write_opt(program& p, program_node* node);
 };
 
+class prepare_primitive_fusing_through : public base_pass {
+public:
+    prepare_primitive_fusing_through() : base_pass("prepare_primitive_fusing_through") {}
+    void run(program& p) override;
+};
+
 class prepare_primitive_fusing : public base_pass {
 public:
     explicit prepare_primitive_fusing(layout_optimizer& lo_ref) :

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

@@ -159,10 +159,10 @@ public:
     program_node& get_dependency(size_t idx) const { return *dependencies.at(idx); }
 
     // replaces idx-th dependency of 'this' with 'new_dep', calls program::remove_if_dangling(old_dep)
-    void replace_dependency(size_t idx, program_node& new_dep);
+    void replace_dependency(size_t idx, program_node& new_dep, bool remove_if_dangling = true);
     // searches for 'old_dep' in dependencies list of 'this' and replaces it with 'new_dep', calls
     // program::remove_if_dangling(old_dep)
-    void replace_dependency(program_node const& old_dep, program_node& new_dep);
+    void replace_dependency(program_node const& old_dep, program_node& new_dep, bool remove_if_dangling = true);
 
     std::vector<primitive_id> get_dependencies_ids() const;
 

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

@@ -509,6 +509,8 @@ void program::pre_optimize_graph(bool is_internal) {
 
     reorder_factory rf;
     if (options.get<build_option_type::optimize_data>()->enabled()) {
+        apply_opt_pass<prepare_primitive_fusing_through>();
+
         apply_opt_pass<pre_replace_deconv>(lo);
 
         apply_opt_pass<prepare_primitive_fusing>(lo);
@@ -877,13 +879,13 @@ void program::swap_names(program_node& node1, program_node& node2) {
     std::swap(_extract_id(node1), _extract_id(node2));
 }
 
-void program::replace_all_usages(program_node& old_node, program_node& new_node) {
+void program::replace_all_usages(program_node& old_node, program_node& new_node, bool remove_if_dangling) {
     // We need a copy of users of old_node because old_node may be removed when doing replace_dependency()
     const std::list<program_node*> users(old_node.users);
     auto itr = users.begin();
     while (itr != users.end()) {
         auto user = *(itr++);
-        user->replace_dependency(old_node, new_node);
+        user->replace_dependency(old_node, new_node, remove_if_dangling);
     }
 }
 
@@ -967,7 +969,7 @@ bool program::remove_if_dangling(program_node& node) {
     return true;
 }
 
-bool program::extract_and_remove(program_node& node) {
+bool program::extract(program_node& node) {
     if (node.get_dependencies().size() != 1)
         return false;
 
@@ -1006,13 +1008,32 @@ bool program::extract_and_remove(program_node& node) {
     node.dependencies.clear();
 
     if (!node.is_endpoint())
-        replace_all_usages(node, input);
-    else
-        remove_if_dangling(node);
+        replace_all_usages(node, input, false);
+
+    if (std::find(processing_order.begin(), processing_order.end(), &node) != processing_order.end())
+        processing_order.erase(&node);
 
     return true;
 }
 
+bool program::extract_and_remove(program_node& node) {
+    if (extract(node)) {
+        return remove_if_dangling(node);
+    }
+
+    return false;
+}
+
+bool program::move_node(program_node& node,
+                        program_node& new_prev,
+                        program_node& new_next) {
+    if (extract(node)) {
+        add_intermediate(node, new_next, new_prev);
+        return true;
+    }
+
+    return false;
+}
 
 void program::fuse_nodes(program_node &fused_node,
                          program_node &peer_node,

src/plugins/intel_gpu/src/graph/program_node.cpp

@@ -32,7 +32,7 @@ program_node::program_node(std::shared_ptr<primitive> prim, program& prog)
         output_layout.data_padding = prim->output_padding;
 }
 
-void program_node::replace_dependency(size_t idx, program_node& new_dep) {
+void program_node::replace_dependency(size_t idx, program_node& new_dep, bool remove_if_dangling) {
     if (idx >= dependencies.size())
         return;
     if (dependencies[idx] == &new_dep)
@@ -47,16 +47,18 @@ void program_node::replace_dependency(size_t idx, program_node& new_dep) {
     if (it != dependencies[idx]->users.end()) {
         dependencies[idx]->users.erase(it);
     }
-    myprog.remove_if_dangling(*dependencies[idx]);
+
+    if (remove_if_dangling)
+        myprog.remove_if_dangling(*dependencies[idx]);
 
     dependencies[idx] = &new_dep;
     new_dep.users.push_back(this);
 }
 
-void program_node::replace_dependency(program_node const& old_dep, program_node& new_dep) {
+void program_node::replace_dependency(program_node const& old_dep, program_node& new_dep, bool remove_if_dangling) {
     for (size_t i = 0; i < dependencies.size(); ++i)
         if (dependencies[i] == &old_dep)
-            return replace_dependency(i, new_dep);
+            return replace_dependency(i, new_dep, remove_if_dangling);
 }
 
 std::vector<primitive_id> program_node::get_dependencies_ids() const {
@@ -104,10 +106,10 @@ std::unique_ptr<json_composite> program_node::desc_to_json() const {
 
     node_info->add("output layout", output_layout_info);
 
-    node_info->add("in data flow", bool_to_str(data_flow));
     node_info->add("constant", bool_to_str(constant));
     node_info->add("in data flow", bool_to_str(data_flow));
     node_info->add("output", bool_to_str(output));
+    node_info->add("optimized", bool_to_str(optimized));
 
     json_composite fused_nodes_info;
     size_t index = 0;

src/plugins/intel_gpu/src/graph/softmax.cpp

@@ -16,7 +16,13 @@ primitive_type_id softmax::type_id() {
 layout softmax_inst::calc_output_layout(softmax_node const& node) {
     assert(static_cast<bool>(node.get_primitive()->output_data_type) == false &&
            "Output data type forcing is not supported for softmax_node!");
-    return node.input().get_output_layout();
+
+    auto output_layout = node.input().get_output_layout();
+
+    if (node.has_fused_primitives())
+        output_layout.data_type = node.get_fused_output_layout().data_type;
+
+    return output_layout;
 }
 
 std::string softmax_inst::to_string(softmax_node const& node) {

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_items_class_kernel_base.cpp

@@ -12,6 +12,8 @@ ParamsKey SoftmaxItemsClassKernelBase::GetDefaultSupportedKey() {
     k.EnableInputDataType(Datatype::F32);
     k.EnableOutputDataType(Datatype::F16);
     k.EnableOutputDataType(Datatype::F32);
+    k.EnableOutputDataType(Datatype::INT8);
+    k.EnableOutputDataType(Datatype::UINT8);
     k.EnableInputLayout(DataLayout::byxf);
     k.EnableInputLayout(DataLayout::bfyx);
     k.EnableInputLayout(DataLayout::yxfb);
@@ -34,6 +36,7 @@ ParamsKey SoftmaxItemsClassKernelBase::GetDefaultSupportedKey() {
     k.EnableSoftmaxDim(SoftmaxDim::Y);
     k.EnableSoftmaxDim(SoftmaxDim::Z);
     k.EnableSoftmaxDim(SoftmaxDim::FEATURE);
+    k.EnableDifferentTypes();
     k.EnableTensorOffset();
     k.EnableTensorPitches();
     k.EnableBatching();
@@ -58,6 +61,8 @@ std::vector<size_t> SoftmaxItemsClassKernelBase::GetSoftmaxDimGlobalSizes(Softma
 JitConstants SoftmaxItemsClassKernelBase::GetJitConstants(const softmax_params& params, DispatchData dispatchData) const {
     auto jit = SoftmaxKernelBase::GetJitConstants(params, dispatchData);
 
+    std::vector<std::string> idx_order;
+    const auto ndims = params.inputs[0].GetDims().size();
     switch (params.dim) {
         case SoftmaxDim::X:
             jit.AddConstants({
@@ -72,6 +77,7 @@ JitConstants SoftmaxItemsClassKernelBase::GetJitConstants(const softmax_params&
                 MakeJitConstant("OUTPUT_OTHER1_PITCH", "OUTPUT_FEATURE_PITCH"),
                 MakeJitConstant("OUTPUT_CLASS_PITCH", "OUTPUT_X_PITCH"),
             });
+            idx_order = {"batch", "other1", ndims == 5 ? "other2" : "0", "other0", "cls"};
             break;
         case SoftmaxDim::Y:
             jit.AddConstants({
@@ -86,6 +92,7 @@ JitConstants SoftmaxItemsClassKernelBase::GetJitConstants(const softmax_params&
                 MakeJitConstant("OUTPUT_OTHER2_PITCH", "OUTPUT_Z_PITCH"),
                 MakeJitConstant("OUTPUT_CLASS_PITCH", "OUTPUT_Y_PITCH"),
             });
+            idx_order = {"batch", "other1", ndims == 5 ? "other2" : "0", "cls", "other0"};
             break;
         case SoftmaxDim::Z:
             jit.AddConstants({
@@ -100,6 +107,7 @@ JitConstants SoftmaxItemsClassKernelBase::GetJitConstants(const softmax_params&
                 MakeJitConstant("OUTPUT_OTHER2_PITCH", "OUTPUT_Y_PITCH"),
                 MakeJitConstant("OUTPUT_CLASS_PITCH", "OUTPUT_Z_PITCH"),
             });
+            idx_order = {"batch", "other1", "cls", "other2", "other0"};
             break;
         case SoftmaxDim::FEATURE:
             jit.AddConstants({
@@ -114,14 +122,21 @@ JitConstants SoftmaxItemsClassKernelBase::GetJitConstants(const softmax_params&
                 MakeJitConstant("OUTPUT_OTHER2_PITCH", "OUTPUT_Z_PITCH"),
                 MakeJitConstant("OUTPUT_CLASS_PITCH", "OUTPUT_FEATURE_PITCH"),
             });
+            idx_order = {"batch", "cls", ndims == 5 ? "other2" : "0", "other1", "other0"};
             break;
         default:
             break;
     }
 
-    // TODO: W/A - currently using low precision accumulator type. (for testing only)
-    if (params.outputs[0].GetDType() == Datatype::F16) {
-        jit.AddConstant(MakeJitConstant("ACCUMULATOR_TYPE", "half"));
+    auto acc_dt = GetAccumulatorType(params);
+    jit.Merge(MakeTypeJitConstants(acc_dt, "ACCUMULATOR"));
+
+    if (!params.fused_ops.empty()) {
+        FusedOpsConfiguration conf = {"",
+                                      idx_order,
+                                      "res",
+                                      acc_dt};
+        jit.Merge(MakeFusedOpsJitConstants(params, {conf}));
     }
 
     return jit;

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_items_class_kernel_base.h

@@ -17,5 +17,14 @@ protected:
     JitConstants GetJitConstants(const softmax_params& params, DispatchData dispatchData) const override;
     static ParamsKey GetDefaultSupportedKey();
     static std::vector<size_t> GetSoftmaxDimGlobalSizes(SoftmaxDim dim, const DataTensor& output);
+    Datatype GetAccumulatorType(const softmax_params& params) const {
+        if (params.inputs[0].GetDType() == Datatype::F16)
+            return Datatype::F16;
+        else
+            return Datatype::F32;
+    }
+    std::vector<KernelBase::FusedOpType> GetSupportedFusedOps() const override {
+        return { FusedOpType::QUANTIZE };
+    }
 };
 }  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_kernel_base.h

@@ -52,6 +52,12 @@ protected:
     virtual JitConstants GetJitConstants(const softmax_params& params, DispatchData dispatchData) const;
     virtual DispatchData SetDefault(const softmax_params& params, const optional_params& optParams) const;
     KernelsData GetCommonKernelsData(const Params& params, const optional_params& optParams) const;
+    Datatype GetActivationType(const softmax_params& params) const {
+        if (params.inputs[0].GetDType() == Datatype::F16)
+            return Datatype::F16;
+        else
+            return Datatype::F32;
+    }
 };
 
 class SoftmaxKernelBaseBF : public SoftmaxKernelBase {

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_kernel_bf.cpp

@@ -12,6 +12,8 @@ ParamsKey SoftmaxKernel_bf::GetSupportedKey() const {
     k.EnableInputDataType(Datatype::F32);
     k.EnableOutputDataType(Datatype::F16);
     k.EnableOutputDataType(Datatype::F32);
+    k.EnableOutputDataType(Datatype::UINT8);
+    k.EnableOutputDataType(Datatype::INT8);
     k.EnableInputLayout(DataLayout::bfyx);
     k.EnableInputLayout(DataLayout::bf);
     k.EnableOutputLayout(DataLayout::bfyx);
@@ -19,6 +21,7 @@ ParamsKey SoftmaxKernel_bf::GetSupportedKey() const {
     k.EnableSoftmaxDim(SoftmaxDim::X);  // in case that it can be flatten
     k.EnableSoftmaxDim(SoftmaxDim::Y);
     k.EnableSoftmaxDim(SoftmaxDim::FEATURE);
+    k.EnableDifferentTypes();
     k.EnableBatching();
     return k;
 }
@@ -64,4 +67,25 @@ KernelsPriority SoftmaxKernel_bf::GetKernelsPriority(const Params& /*params*/, c
 KernelsData SoftmaxKernel_bf::GetKernelsData(const Params& params, const optional_params& optionalParams) const {
     return GetCommonKernelsData(params, optionalParams);
 }
-}  // namespace kernel_selector
+
+JitConstants SoftmaxKernel_bf::GetJitConstants(const softmax_params& params, DispatchData dispatchData) const {
+    auto jit = Parent::GetJitConstants(params, dispatchData);
+
+    jit.Merge(MakeTypeJitConstants(GetActivationType(params), "ACTIVATION"));
+
+    if (!params.fused_ops.empty()) {
+        auto input_dt = GetActivationType(params);
+        FusedOpsConfiguration conf_main = {"_MAIN",
+                                           {"data_set_offset", "in_data_set_idx + i * workers_per_data_set", "0", "0"},
+                                           "dequantized",
+                                           input_dt};
+        FusedOpsConfiguration conf_leftovers = {"_LEFTOVERS",
+                                                {"data_set_offset", "workers_per_data_set * ITEMS_NUM + in_data_set_idx", "0", "0"},
+                                                "dequantized",
+                                                input_dt};
+        jit.Merge(MakeFusedOpsJitConstants(params, {conf_main, conf_leftovers}));
+    }
+
+    return jit;
+}
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_kernel_bf.h

@@ -16,6 +16,12 @@ public:
     KernelsData GetKernelsData(const Params& params, const optional_params& options) const override;
     KernelsPriority GetKernelsPriority(const Params& params, const optional_params& options) const override;
     ParamsKey GetSupportedKey() const override;
+
+protected:
     DispatchData SetDefault(const softmax_params& params, const optional_params& optParams) const override;
+    JitConstants GetJitConstants(const softmax_params& params, DispatchData dispatchData) const override;
+    std::vector<KernelBase::FusedOpType> GetSupportedFusedOps() const override {
+        return { FusedOpType::QUANTIZE };
+    }
 };
-}  // namespace kernel_selector
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_kernel_fb.cpp

@@ -95,4 +95,25 @@ KernelsData SoftmaxKernel_fb::GetKernelsData(const Params& params, const optiona
     }
     return GetCommonKernelsData(params, optParams);
 }
-}  // namespace kernel_selector
+
+JitConstants SoftmaxKernel_fb::GetJitConstants(const softmax_params& params, DispatchData dispatchData) const {
+    auto jit = Parent::GetJitConstants(params, dispatchData);
+
+    jit.Merge(MakeTypeJitConstants(GetActivationType(params), "ACTIVATION"));
+
+    if (!params.fused_ops.empty()) {
+        auto input_dt = GetActivationType(params);
+        FusedOpsConfiguration conf_main = {"_MAIN",
+                                           {"global_id", "LWS * i", "0", "0"},
+                                           "dequantized",
+                                           input_dt};
+        FusedOpsConfiguration conf_leftovers = {"_LEFTOVERS",
+                                                {"global_id", "LWS * ITEMS_NUM", "0", "0"},
+                                                "dequantized",
+                                                input_dt};
+        jit.Merge(MakeFusedOpsJitConstants(params, {conf_main, conf_leftovers}));
+    }
+
+    return jit;
+}
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/softmax/softmax_kernel_fb.h

@@ -20,5 +20,9 @@ public:
 protected:
     bool Validate(const Params& p, const optional_params& o) const override;
     DispatchData SetDefault(const softmax_params& params, const optional_params& optParams) const override;
+    JitConstants GetJitConstants(const softmax_params& params, DispatchData dispatchData) const override;
+    std::vector<KernelBase::FusedOpType> GetSupportedFusedOps() const override {
+        return { FusedOpType::QUANTIZE };
+    }
 };
-}  // namespace kernel_selector
+}  // namespace kernel_selector

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

@@ -7,23 +7,28 @@
 
 
 __attribute__((reqd_work_group_size(LWS, 1, 1)))
-KERNEL (softmax_gpu_continoues_bfyx)(const __global UNIT_TYPE* input, __global UNIT_TYPE* output)
-{
+KERNEL (softmax_gpu_continuous_bfyx)(
+    const __global INPUT0_TYPE* input,
+    __global OUTPUT_TYPE* output
+#if HAS_FUSED_OPS_DECLS
+    , FUSED_OPS_DECLS
+#endif
+) {
     const uint data_set_idx = get_global_id(1);     //in processing of which data set this WI participates?
     const uint workers_per_data_set = LWS;          //how many WI participates in processing of one data set
     const uint in_data_set_idx = get_global_id(0);  //this WI's id in group of items processing single data set
     const uint data_set_size = DATA_SET_SIZE;       //how many elements are in one data set
-    const uint data_sets_count = DATA_SETS_COUNT;   //how many data sets are in the processing payload     
+    const uint data_sets_count = DATA_SETS_COUNT;   //how many data sets are in the processing payload
 
     const uint data_set_offset = data_set_idx * data_set_size;
     const uint my_data_offset = data_set_offset + in_data_set_idx;
 
-    UNIT_TYPE my_chunk[ITEMS_NUM + 1];
-    UNIT_TYPE my_maximum = -UNIT_VAL_MAX;
-    UNIT_TYPE my_sum = UNIT_VAL_ZERO;
-    UNIT_TYPE tmp;
+    INPUT0_TYPE my_chunk[ITEMS_NUM + 1];
+    INPUT0_TYPE my_maximum = -UNIT_VAL_MAX;
+    INPUT0_TYPE my_sum = UNIT_VAL_ZERO;
+    INPUT0_TYPE tmp;
 
-    __local UNIT_TYPE lg_storage[LWS];
+    __local INPUT0_TYPE lg_storage[LWS];
 
     //each WI reads ITEMS_NUM consecutive items from batch
     for (uint i=0; i<ITEMS_NUM; ++i)
@@ -85,8 +90,23 @@ KERNEL (softmax_gpu_continoues_bfyx)(const __global UNIT_TYPE* input, __global U
 
     my_sum = lg_storage[0];
 
+#if HAS_FUSED_OPS
+    for (uint i=0; i<ITEMS_NUM; ++i)
+    {
+        ACTIVATION_TYPE dequantized = my_chunk[i] / my_sum;
+        FUSED_OPS_MAIN;
+        output[my_data_offset + i * workers_per_data_set] = FUSED_OPS_RESULT_MAIN;
+    }
+    if (in_data_set_idx < LEFTOVERS)
+    {
+        ACTIVATION_TYPE dequantized = my_chunk[ITEMS_NUM] / my_sum;
+        FUSED_OPS_LEFTOVERS;
+        output[data_set_offset + workers_per_data_set * ITEMS_NUM + in_data_set_idx] = FUSED_OPS_RESULT_LEFTOVERS;
+    }
+#else
     for (uint i=0; i<ITEMS_NUM; ++i)
         output[my_data_offset + i * workers_per_data_set] = ACTIVATION(my_chunk[i] / my_sum, ACTIVATION_PARAMS);
     if (in_data_set_idx < LEFTOVERS)
         output[data_set_offset + workers_per_data_set * ITEMS_NUM + in_data_set_idx] = ACTIVATION(my_chunk[ITEMS_NUM] / my_sum, ACTIVATION_PARAMS);
+#endif
 }

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

@@ -6,9 +6,9 @@
 #include "include/batch_headers/data_types.cl"
 
 
-UNIT_TYPE FUNC(find_max_value)(__local UNIT_TYPE* partial_max, const int global_id, const int idx, const int batch_offset, const int data_sets_count, const __global UNIT_TYPE* input)
+INPUT0_TYPE FUNC(find_max_value)(__local INPUT0_TYPE* partial_max, const int global_id, const int idx, const int batch_offset, const int data_sets_count, const __global INPUT0_TYPE* input)
 {
-    UNIT_TYPE value = -UNIT_VAL_MAX;
+    INPUT0_TYPE value = -UNIT_VAL_MAX;
     for(int i = 0; i < ITEMS_NUM; i++)
     {
         value = max(value, input[LWS * i + global_id]);
@@ -28,19 +28,24 @@ UNIT_TYPE FUNC(find_max_value)(__local UNIT_TYPE* partial_max, const int global_
     return partial_max[batch_offset];
 }
 
-KERNEL (softmax_gpu_continoues_yxfb)(const __global UNIT_TYPE* input, __global UNIT_TYPE* output)
-{
-    const uint data_sets_count = DATA_SETS_COUNT;   //how many data sets are in the processing payload     
+KERNEL (softmax_gpu_continuous_yxfb)(
+    const __global UNIT_TYPE* input,
+    __global UNIT_TYPE* output
+#if HAS_FUSED_OPS_DECLS
+    , FUSED_OPS_DECLS
+#endif
+) {
+    const uint data_sets_count = DATA_SETS_COUNT;   //how many data sets are in the processing payload
 
     const int global_id = get_global_id(0);
     const int idx = global_id / data_sets_count;
 
     const int batch_offset = global_id % data_sets_count;
 
-    __local UNIT_TYPE partial_max[LWS];
-    const UNIT_TYPE max_value = FUNC_CALL(find_max_value)(partial_max, global_id, idx, batch_offset, data_sets_count, input);
+    __local INPUT0_TYPE partial_max[LWS];
+    const INPUT0_TYPE max_value = FUNC_CALL(find_max_value)(partial_max, global_id, idx, batch_offset, data_sets_count, input);
 
-    UNIT_TYPE tmp_vals[ITEMS_NUM + 1];
+    INPUT0_TYPE tmp_vals[ITEMS_NUM + 1];
     for(int i = 0; i < ITEMS_NUM; i++)
     {
         tmp_vals[i] = native_exp(input[LWS * i + global_id] - max_value);
@@ -48,7 +53,7 @@ KERNEL (softmax_gpu_continoues_yxfb)(const __global UNIT_TYPE* input, __global U
     tmp_vals[ITEMS_NUM] = global_id < LEFTOVERS ? native_exp(input[LWS * ITEMS_NUM + global_id] - max_value) : UNIT_VAL_ZERO;
 
     // accumulate all values;
-    __local UNIT_TYPE partial_acc[LWS]; // all values accumulated;
+    __local INPUT0_TYPE partial_acc[LWS]; // all values accumulated;
     partial_acc[global_id] = UNIT_VAL_ZERO;
     for(int i = 0; i < ITEMS_NUM + 1; i++)
     {
@@ -64,10 +69,25 @@ KERNEL (softmax_gpu_continoues_yxfb)(const __global UNIT_TYPE* input, __global U
         }
     }
     barrier(CLK_LOCAL_MEM_FENCE);
+#if HAS_FUSED_OPS
+    for(int i = 0; i < ITEMS_NUM; i++)
+    {
+        ACTIVATION_TYPE dequantized = tmp_vals[i] / partial_acc[batch_offset];
+        FUSED_OPS_MAIN;
+        output[LWS * i + global_id] = FUSED_OPS_RESULT_MAIN;
+    }
+    if(global_id < LEFTOVERS)
+    {
+        ACTIVATION_TYPE dequantized = tmp_vals[ITEMS_NUM] / partial_acc[batch_offset];
+        FUSED_OPS_LEFTOVERS;
+        output[LWS * ITEMS_NUM + global_id] = FUSED_OPS_RESULT_LEFTOVERS;
+    }
+#else
     for(int i = 0; i < ITEMS_NUM; i++)
     {
         output[LWS * i + global_id] = ACTIVATION(tmp_vals[i] / partial_acc[batch_offset], ACTIVATION_PARAMS);
     }
     if(global_id < LEFTOVERS)
         output[LWS * ITEMS_NUM + global_id] = ACTIVATION(tmp_vals[ITEMS_NUM] / partial_acc[batch_offset], ACTIVATION_PARAMS);
-}
+#endif
+}

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

@@ -3,15 +3,19 @@
 //
 
 #include "include/batch_headers/common.cl"
-#include "include/acc_type.cl"
 
 
 #define DATA_PER_WORKITEM ( (INPUT0_CLASS_NUM + (WORKITEMS_PER_CLASSES - 1) ) / WORKITEMS_PER_CLASSES)
 #define FULL_ITERATIONS_NUM (INPUT0_CLASS_NUM / WORKITEMS_PER_CLASSES)
 
 __attribute__((intel_reqd_sub_group_size(16)))
-KERNEL(softmax_items_class_optimized)(__global INPUT0_TYPE* input, __global OUTPUT_TYPE* output)
-{
+KERNEL(softmax_items_class_optimized)(
+    __global INPUT0_TYPE* input,
+    __global OUTPUT_TYPE* output
+#if HAS_FUSED_OPS_DECLS
+    , FUSED_OPS_DECLS
+#endif
+) {
 #if INPUT0_DIMS == 5
     const uint other0 = (uint)get_group_id(0) % INPUT0_OTHER0_SIZE;
     const uint other2 = (uint)get_group_id(0) / INPUT0_OTHER0_SIZE;
@@ -26,21 +30,21 @@ KERNEL(softmax_items_class_optimized)(__global INPUT0_TYPE* input, __global OUTP
     const uint in_depth_offset  = batch*INPUT0_BATCH_PITCH + other2*INPUT0_OTHER2_PITCH + other1*INPUT0_OTHER1_PITCH + other0*INPUT0_OTHER0_PITCH + INPUT0_OFFSET;
     const uint out_depth_offset = batch*OUTPUT_BATCH_PITCH + other2*OUTPUT_OTHER2_PITCH + other1*OUTPUT_OTHER1_PITCH + other0*OUTPUT_OTHER0_PITCH + OUTPUT_OFFSET;
 
-    UNIT_TYPE max_value = UNIT_VAL_MIN;
-    UNIT_TYPE data[DATA_PER_WORKITEM];
+    ACCUMULATOR_TYPE max_value = UNIT_VAL_MIN;
+    ACCUMULATOR_TYPE data[DATA_PER_WORKITEM];
 
     // PART 1. Calculate MAX value
     uint input_idx = in_depth_offset + simd_lane * INPUT0_CLASS_PITCH;
     for (uint cls = 0; cls < FULL_ITERATIONS_NUM; cls++)
     {
-        UNIT_TYPE in = input[input_idx];
+        ACCUMULATOR_TYPE in = input[input_idx];
         max_value = max(max_value, in);
         data[cls] = in;
         input_idx += WORKITEMS_PER_CLASSES*INPUT0_CLASS_PITCH;
     }
     if(simd_lane < LEFTOVERS)
     {
-        UNIT_TYPE in = input[input_idx];
+        ACCUMULATOR_TYPE in = input[input_idx];
         max_value = max(max_value, in);
         data[DATA_PER_WORKITEM-1] = in;
     }
@@ -51,8 +55,8 @@ KERNEL(softmax_items_class_optimized)(__global INPUT0_TYPE* input, __global OUTP
     ACCUMULATOR_TYPE denominator = 0.0;
     for (uint cls = 0; cls < FULL_ITERATIONS_NUM; cls++)
     {
-// This is a temporary solution for unresolved problem when ocl kernels compilation step doesn't produce actual binaries 
-// for current kernel but driver doesn't report any errors (JIRA CVS-32211)
+// This is a temporary solution for unresolved problem when ocl kernels compilation step doesn't produce actual binaries
+// for current kernel but driver doesn't report any errors (JIRA 32211)
 #if HAS_DRIVER_PROBLEMS
         data[cls] = data[cls] == max_value ? 1.0 : native_exp(data[cls] - max_value);
 #else
@@ -62,8 +66,8 @@ KERNEL(softmax_items_class_optimized)(__global INPUT0_TYPE* input, __global OUTP
     }
     if(simd_lane < LEFTOVERS)
     {
-// This is a temporary solution for unresolved problem when ocl kernels compilation step doesn't produce actual binaries 
-// for current kernel but driver doesn't report any errors (JIRA CVS-32211)
+// This is a temporary solution for unresolved problem when ocl kernels compilation step doesn't produce actual binaries
+// for current kernel but driver doesn't report any errors (JIRA 32211)
 #if HAS_DRIVER_PROBLEMS
         data[DATA_PER_WORKITEM-1] = data[DATA_PER_WORKITEM-1] == max_value ? 1.0 : native_exp(data[DATA_PER_WORKITEM-1] - max_value);
 #else
@@ -78,14 +82,24 @@ KERNEL(softmax_items_class_optimized)(__global INPUT0_TYPE* input, __global OUTP
     uint output_idx = out_depth_offset + simd_lane * OUTPUT_CLASS_PITCH;
     for (uint cls = 0; cls < FULL_ITERATIONS_NUM; cls++)
     {
-        const UNIT_TYPE res = data[cls] / (UNIT_TYPE)denominator;
+        const ACCUMULATOR_TYPE res = data[cls] / denominator;
+#if HAS_FUSED_OPS
+        FUSED_OPS;
+        output[output_idx] = FUSED_OPS_RESULT;
+#else
         output[output_idx] = ACTIVATION(res, ACTIVATION_PARAMS);
+#endif
         output_idx += WORKITEMS_PER_CLASSES * OUTPUT_CLASS_PITCH;
     }
     if(simd_lane < LEFTOVERS)
     {
-        const UNIT_TYPE res = data[DATA_PER_WORKITEM-1] / (UNIT_TYPE)denominator;
+        const ACCUMULATOR_TYPE res = data[DATA_PER_WORKITEM-1] / denominator;
+#if HAS_FUSED_OPS
+        FUSED_OPS;
+        output[output_idx] = FUSED_OPS_RESULT;
+#else
         output[output_idx] = ACTIVATION(res, ACTIVATION_PARAMS);
+#endif
     }
 }
 

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

@@ -2,13 +2,17 @@
 // SPDX-License-Identifier: Apache-2.0
 //
 
-#include "include/acc_type.cl"
 #include "include/batch_headers/common.cl"
 
 
 __attribute__((intel_reqd_sub_group_size(16)))
-KERNEL(softmax)(__global INPUT0_TYPE* input, __global OUTPUT_TYPE* output)
-{
+KERNEL(softmax)(
+    __global INPUT0_TYPE* input,
+    __global OUTPUT_TYPE* output
+#if HAS_FUSED_OPS_DECLS
+    , FUSED_OPS_DECLS
+#endif
+) {
 #if INPUT0_DIMS == 5
     const uint other0 = (uint)get_global_id(0) % INPUT0_OTHER0_SIZE;
     const uint other2 = (uint)get_global_id(0) / INPUT0_OTHER0_SIZE;
@@ -22,13 +26,13 @@ KERNEL(softmax)(__global INPUT0_TYPE* input, __global OUTPUT_TYPE* output)
     const uint in_depth_offset  = batch*INPUT0_BATCH_PITCH + other2*INPUT0_OTHER2_PITCH + other1*INPUT0_OTHER1_PITCH + other0*INPUT0_OTHER0_PITCH + INPUT0_OFFSET;
     const uint out_depth_offset = batch*OUTPUT_BATCH_PITCH + other2*OUTPUT_OTHER2_PITCH + other1*OUTPUT_OTHER1_PITCH + other0*OUTPUT_OTHER0_PITCH + OUTPUT_OFFSET;
 
-    UNIT_TYPE max_value = UNIT_VAL_MIN;
-    UNIT_TYPE data[INPUT0_CLASS_NUM];
+    ACCUMULATOR_TYPE max_value = UNIT_VAL_MIN;
+    ACCUMULATOR_TYPE data[INPUT0_CLASS_NUM];
 
     for (uint cls = 0; cls < INPUT0_CLASS_NUM; ++cls)
     {
         const uint index = in_depth_offset + cls*INPUT0_CLASS_PITCH;
-        UNIT_TYPE in = input[index];
+        ACCUMULATOR_TYPE in = input[index];
         max_value = max(max_value, in);
         data[cls] = in;
     }
@@ -43,8 +47,13 @@ KERNEL(softmax)(__global INPUT0_TYPE* input, __global OUTPUT_TYPE* output)
 
     for (uint cls = 0; cls < INPUT0_CLASS_NUM; ++cls)
     {
-        const UNIT_TYPE res = data[cls] / (UNIT_TYPE)denominator;
+        const ACCUMULATOR_TYPE res = data[cls] / denominator;
         const uint output_idx = out_depth_offset + cls*OUTPUT_CLASS_PITCH;
+#if HAS_FUSED_OPS
+        FUSED_OPS;
+        output[output_idx] = FUSED_OPS_RESULT;
+#else
         output[output_idx] = ACTIVATION(res, ACTIVATION_PARAMS);
+#endif
     }
 }

src/plugins/intel_gpu/src/kernel_selector/core/common/jitter.cpp

@@ -1532,7 +1532,6 @@ JitConstants FusedOpsCodeGenerator::MakeOpJitConstants(const FusedOpsConfigurati
 
     std::string op_decls = "";
     auto vec_size = conf.vec_size;
-    auto idx = conf.bfzyx_idx_order;
     std::string shuffle_var = conf.shuffle_var_name;
     bool is_shuffled = false;
 

src/plugins/intel_gpu/tests/fusions/convolution_fusion_test.cpp

@@ -115,7 +115,7 @@ public:
 
 class ConvReorderFusingTest : public BaseFusingTest<convolution_test_params> {
 public:
-    void execute(convolution_test_params& p) {
+    void execute(convolution_test_params& p, std::map<std::string, std::vector<std::string>> expected_fused_primitives_ids = {}) {
         auto input_prim = get_mem(get_input_layout(p));
         network network_not_fused(this->engine, this->topology_non_fused, bo_not_fused);
         network network_fused(this->engine, this->topology_fused, bo_fused);
@@ -123,6 +123,7 @@ public:
         network_not_fused.set_input_data("input", input_prim);
 
         compare(network_not_fused, network_fused, p, true);
+        check_fusions_correctness(network_fused, expected_fused_primitives_ids);
     }
 
     layout get_input_layout(convolution_test_params& p) {
@@ -2641,6 +2642,38 @@ TEST_P(conv_fp32_reorder_bfyx_to_fsv32_conv_fused_activation, have_fused_activat
     auto dw_weights_layout = layout{ p.default_type, format::goiyx, dw_tensor };
     ov::Strides dw_stride = { 1, 1 };
 
+    create_topologies(
+        input_layout("input", get_input_layout(p)),
+        data("weights", get_mem(get_weights_layout(p), -127, 127)),
+        data("weights_dw", get_mem(dw_weights_layout, -127, 127)),
+        data("actv_params", get_mem(get_per_channel_layout(p), -127, 127)),
+        convolution("conv_prim", "input", { "weights" }, p.groups, p.stride, p.pad, p.dilation),
+        reorder("reorder_fsv32", "conv_prim", format::fs_b_yx_fsv32, data_types::f32),
+        activation("activation_quantize", "reorder_fsv32", "actv_params", activation_func::relu_negative_slope),
+        convolution("conv_prim2", "activation_quantize", { "weights_dw" }, p.out_shape.feature[0], dw_stride, p.pad, p.dilation),
+        activation("activation", "conv_prim2", activation_func::abs)
+    );
+
+    implementation_desc conv_impl = { format::fs_b_yx_fsv32, "" };
+    bo_fused.set_option(build_option::force_implementations({ { "conv_prim2", conv_impl } }));
+    bo_fused.set_option(build_option::force_implementations({ { "activation", conv_impl } }));
+
+    execute(p);
+}
+
+INSTANTIATE_TEST_SUITE_P(fusings_gpu, conv_fp32_reorder_bfyx_to_fsv32_conv_fused_activation, ::testing::ValuesIn(std::vector<convolution_test_params>{
+    convolution_test_params{ FSV32_CASE_CONV_FP32_1, 5, 6 }
+}));
+
+// activation will be fused through 'reorder_fsv32' and 'reorder_fsv32' will be fused as well
+class conv_fp32_reorder_bfyx_to_fsv32_conv_fused_through_activation : public ConvReorderFusingTest {};
+TEST_P(conv_fp32_reorder_bfyx_to_fsv32_conv_fused_through_activation, have_fused_through_activation) {
+    auto p = GetParam();
+
+    auto dw_tensor = cldnn::tensor(group(p.out_shape.feature[0]), batch(1), feature(1), spatial(2, 2));
+    auto dw_weights_layout = layout{ p.default_type, format::goiyx, dw_tensor };
+    ov::Strides dw_stride = { 1, 1 };
+
     create_topologies(
         input_layout("input", get_input_layout(p)),
         data("weights", get_mem(get_weights_layout(p), -127, 127)),
@@ -2656,10 +2689,11 @@ TEST_P(conv_fp32_reorder_bfyx_to_fsv32_conv_fused_activation, have_fused_activat
     bo_fused.set_option(build_option::force_implementations({ { "conv_prim2", conv_impl } }));
     bo_fused.set_option(build_option::force_implementations({ { "activation", conv_impl } }));
 
-    execute(p);
+    execute(p, {{"conv_prim", {"activation_quantize"}}});
 }
-INSTANTIATE_TEST_SUITE_P(fusings_gpu, conv_fp32_reorder_bfyx_to_fsv32_conv_fused_activation, ::testing::ValuesIn(std::vector<convolution_test_params>{
-    convolution_test_params{ FSV32_CASE_CONV_FP32_1,  5, 6 }
+
+INSTANTIATE_TEST_SUITE_P(fusings_gpu, conv_fp32_reorder_bfyx_to_fsv32_conv_fused_through_activation, ::testing::ValuesIn(std::vector<convolution_test_params>{
+    convolution_test_params{ FSV32_CASE_CONV_FP32_1, 4, 6 }
 }));
 
 // 'reorder_fsv32' is being fused even if it has 'padding'

src/plugins/intel_gpu/tests/fusions/fusion_test_common.hpp

@@ -94,6 +94,24 @@ public:
         }
     }
 
+    void check_fusions_correctness(network& network_fused, std::map<std::string, std::vector<std::string>> expected_fused_primitives_ids = {}) {
+        if (expected_fused_primitives_ids.size()) {
+            auto primitives_info = network_fused.get_primitives_info();
+            for (auto& prim : expected_fused_primitives_ids) {
+                auto info = std::find_if(primitives_info.begin(), primitives_info.end(),
+                                         [&prim](const primitive_info& info) -> bool { return info.original_id == prim.first; });
+                if (info != primitives_info.end()) {
+                    auto fused_primitives = info->c_fused_ids;
+                    for (auto& expected_fused_prim : prim.second)
+                        if (std::find(fused_primitives.begin(), fused_primitives.end(), expected_fused_prim) == fused_primitives.end())
+                            FAIL() << "Couldn't find requested fused primitive id " + prim.first;
+                } else {
+                    FAIL() << "Couldn't find requested primitive id " + prim.first;
+                }
+            }
+        }
+    }
+
     cldnn::memory::ptr get_mem(cldnn::layout l) {
         auto prim = engine.allocate_memory(l);
         tensor s = l.size;

src/plugins/intel_gpu/tests/fusions/resample_fusion_test.cpp

@@ -10,6 +10,7 @@
 #include <intel_gpu/primitives/eltwise.hpp>
 #include <intel_gpu/primitives/data.hpp>
 #include <intel_gpu/primitives/resample.hpp>
+#include <intel_gpu/primitives/reshape.hpp>
 
 #include <cmath>
 
@@ -32,7 +33,7 @@ struct resample_test_params {
 class ResamplePrimitiveFusingTest : public ::BaseFusingTest<resample_test_params> {
 public:
 
-    void execute(resample_test_params& p) {
+    void execute(resample_test_params& p, std::map<std::string, std::vector<std::string>> expected_fused_primitives_ids = {}) {
         auto input_prim = get_mem(get_input_layout(p));
         network network_not_fused(this->engine, this->topology_non_fused, bo_not_fused);
         network network_fused(this->engine, this->topology_fused, bo_fused);
@@ -40,6 +41,7 @@ public:
         network_not_fused.set_input_data("input", input_prim);
 
         compare(network_not_fused, network_fused, p);
+        check_fusions_correctness(network_fused, expected_fused_primitives_ids);
     }
 
     layout get_input_layout(resample_test_params& p) {
@@ -308,3 +310,82 @@ INSTANTIATE_TEST_SUITE_P(fusings_gpu, resample_scale_concat, ::testing::ValuesIn
     resample_test_params{ CASE_RESAMPLE_U8_3, 3, 6 },
     resample_test_params{ CASE_RESAMPLE_U8_4, 3, 6 },
 }));
+
+class resample_scale_fusing_through : public ResamplePrimitiveFusingTest {};
+TEST_P(resample_scale_fusing_through, reshape) {
+    auto p = GetParam();
+    auto reshape_shape = p.out_shape;
+    reshape_shape.feature[0] *= reshape_shape.spatial[0];
+    reshape_shape.spatial[0] = 1;
+
+    create_topologies(
+        input_layout("input", get_input_layout(p)),
+        data("scale_data", get_mem(layout{ p.default_type, p.default_format, tensor{ 1, 1, 1, 1 } })),
+        resample("resample_prim", "input", p.out_shape, p.in_shape.feature[0], p.type),
+        reshape("reshape", "resample_prim", reshape_shape),
+        eltwise("scale", "reshape", "scale_data", eltwise_mode::prod),
+        reorder("reorder_bfyx", "scale", p.default_format, data_types::f32)
+    );
+
+    tolerance = 1e-5f;
+    execute(p, {{"resample_prim", {"scale"}}});
+}
+
+INSTANTIATE_TEST_SUITE_P(fusings_gpu, resample_scale_fusing_through, ::testing::ValuesIn(std::vector<resample_test_params>{
+    resample_test_params{ CASE_RESAMPLE_FP32_1, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP32_2, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP32_3, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP32_4, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP32_5, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP32_6, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP32_7, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP32_8, 2, 3 },
+
+    resample_test_params{ CASE_RESAMPLE_FP16_1, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP16_2, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP16_3, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP16_4, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP16_5, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP16_6, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP16_7, 2, 3 },
+    resample_test_params{ CASE_RESAMPLE_FP16_8, 2, 3 },
+
+    resample_test_params{ CASE_RESAMPLE_I8_1, 2, 4 },
+    resample_test_params{ CASE_RESAMPLE_I8_2, 2, 4 },
+    resample_test_params{ CASE_RESAMPLE_I8_3, 2, 4 },
+    resample_test_params{ CASE_RESAMPLE_I8_4, 2, 4 },
+
+    resample_test_params{ CASE_RESAMPLE_U8_1, 2, 4 },
+    resample_test_params{ CASE_RESAMPLE_U8_2, 2, 4 },
+    resample_test_params{ CASE_RESAMPLE_U8_3, 2, 4 },
+    resample_test_params{ CASE_RESAMPLE_U8_4, 2, 4 },
+}));
+
+class resample_scale_fusing_through_not_allowed : public ResamplePrimitiveFusingTest {};
+TEST_P(resample_scale_fusing_through_not_allowed, reshape_two_users) {
+    auto p = GetParam();
+    auto reshape_shape = p.out_shape;
+    reshape_shape.feature[0] *= reshape_shape.spatial[0];
+    reshape_shape.spatial[0] = 1;
+
+    create_topologies(
+        input_layout("input", get_input_layout(p)),
+        data("scale_data", get_mem(layout{ p.default_type, p.default_format, tensor{ 1, 1, 1, 1 } })),
+        resample("resample_prim", "input", p.out_shape, p.in_shape.feature[0], p.type),
+        reshape("reshape", "resample_prim", reshape_shape),
+        eltwise("scale", "reshape", "scale_data", eltwise_mode::prod),
+        eltwise("sum", "reshape", "scale", eltwise_mode::sum),
+        reorder("reorder_bfyx", "sum", p.default_format, data_types::f32)
+    );
+
+    tolerance = 1e-5f;
+    execute(p);
+}
+
+INSTANTIATE_TEST_SUITE_P(fusings_gpu, resample_scale_fusing_through_not_allowed, ::testing::ValuesIn(std::vector<resample_test_params>{
+    resample_test_params{ CASE_RESAMPLE_FP32_1, 4, 4 },
+    resample_test_params{ CASE_RESAMPLE_FP32_2, 4, 4 },
+
+    resample_test_params{ CASE_RESAMPLE_FP16_1, 4, 4 },
+    resample_test_params{ CASE_RESAMPLE_FP16_2, 4, 4 },
+}));

src/plugins/intel_gpu/tests/fusions/softmax_fusion_test.cpp

@@ -0,0 +1,171 @@
+// Copyright (C) 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include "test_utils.h"
+#include "fusion_test_common.hpp"
+
+#include <intel_gpu/primitives/input_layout.hpp>
+#include <intel_gpu/primitives/quantize.hpp>
+#include <intel_gpu/primitives/eltwise.hpp>
+#include <intel_gpu/primitives/data.hpp>
+#include <intel_gpu/primitives/reshape.hpp>
+#include <intel_gpu/primitives/reorder.hpp>
+
+#include <cmath>
+
+using namespace cldnn;
+using namespace ::tests;
+
+namespace {
+struct softmax_test_params {
+    tensor in_shape;
+    data_types data_type;
+    format input_format;
+    softmax::dimension_t dimension;
+    data_types default_type;
+    format default_format;
+    size_t expected_fused_primitives;
+    size_t expected_not_fused_primitives;
+};
+
+class SoftmaxPrimitiveFusingTest : public ::BaseFusingTest<softmax_test_params> {
+public:
+
+    void execute(softmax_test_params& p, std::map<std::string, std::vector<std::string>> expected_fused_primitives_ids = {}) {
+        auto input_prim = get_mem(get_input_layout(p));
+        network network_not_fused(this->engine, this->topology_non_fused, bo_not_fused);
+        network network_fused(this->engine, this->topology_fused, bo_fused);
+        network_fused.set_input_data("input", input_prim);
+        network_not_fused.set_input_data("input", input_prim);
+
+        compare(network_not_fused, network_fused, p);
+        check_fusions_correctness(network_fused, expected_fused_primitives_ids);
+    }
+
+    layout get_input_layout(softmax_test_params& p) {
+        return layout{ p.data_type, p.input_format, p.in_shape, padding{} };
+    }
+
+    tensor get_reshape_shape(softmax_test_params& p) {
+        auto output_shape = p.in_shape;
+        output_shape.feature[0] *= output_shape.spatial[0];
+        output_shape.spatial[0] = 1;
+        return output_shape;
+    }
+
+    layout get_output_layout(softmax_test_params& p) {
+        return { data_types::f32, p.default_format, get_reshape_shape(p) };
+    }
+};
+
+}  // namespace
+
+
+/* ----------------------------------------------------------------------------------------------------- */
+/* ---------------------------------------- SoftMax cases ---------------------------------------------- */
+/* ----------------------------------------------------------------------------------------------------- */
+#define CASE_SOFTMAX_FP32_1 {1, 15, 4, 5}, data_types::f32, format::bfyx, softmax::dimension_t::normalize_f, data_types::f32, format::bfyx
+#define CASE_SOFTMAX_FP32_2 {1, 15, 4, 5}, data_types::f32, format::bfyx, softmax::dimension_t::normalize_x, data_types::f32, format::bfyx
+#define CASE_SOFTMAX_FP32_3 {1, 15, 4, 5}, data_types::f32, format::bfyx, softmax::dimension_t::normalize_fyx, data_types::f32, format::bfyx
+
+#define CASE_SOFTMAX_FP16_1 {1, 15, 4, 5}, data_types::f16, format::bfyx, softmax::dimension_t::normalize_f, data_types::f16, format::bfyx
+#define CASE_SOFTMAX_FP16_2 {1, 15, 4, 5}, data_types::f16, format::bfyx, softmax::dimension_t::normalize_x, data_types::f16, format::bfyx
+#define CASE_SOFTMAX_FP16_3 {1, 15, 4, 5}, data_types::f16, format::bfyx, softmax::dimension_t::normalize_fyx, data_types::f16, format::bfyx
+
+class softmax_quantize : public SoftmaxPrimitiveFusingTest {};
+TEST_P(softmax_quantize, basic) {
+    auto p = GetParam();
+    create_topologies(input_layout("input", get_input_layout(p)),
+        data("in_lo", get_mem(get_single_element_layout(p), min_random, 0)),
+        data("in_hi", get_mem(get_single_element_layout(p), 1, max_random)),
+        data("out_lo", get_mem(get_single_element_layout(p), -127)),
+        data("out_hi", get_mem(get_single_element_layout(p), 127)),
+        softmax("softmax", "input", p.dimension),
+        quantize("quantize", "softmax", "in_lo", "in_hi", "out_lo", "out_hi", 255, data_types::i8),
+        reorder("reorder_bfyx", "quantize", get_output_layout(p))
+    );
+
+    tolerance = 1.f;
+    execute(p);
+}
+
+INSTANTIATE_TEST_SUITE_P(fusings_gpu, softmax_quantize,
+    ::testing::ValuesIn(std::vector<softmax_test_params>{
+                        softmax_test_params{ CASE_SOFTMAX_FP32_1, 2, 3 },
+                        softmax_test_params{ CASE_SOFTMAX_FP32_2, 3, 3 },
+                        softmax_test_params{ CASE_SOFTMAX_FP32_3, 3, 3 },
+
+                        softmax_test_params{ CASE_SOFTMAX_FP16_1, 2, 3 },
+                        softmax_test_params{ CASE_SOFTMAX_FP16_2, 3, 3 },
+                        softmax_test_params{ CASE_SOFTMAX_FP16_3, 3, 3 },
+}));
+
+class softmax_quantize_fusing_through : public SoftmaxPrimitiveFusingTest {};
+TEST_P(softmax_quantize_fusing_through, reshape) {
+    auto p = GetParam();
+    create_topologies(input_layout("input", get_input_layout(p)),
+        data("in_lo", get_mem(get_single_element_layout(p), min_random, 0)),
+        data("in_hi", get_mem(get_single_element_layout(p), 1, max_random)),
+        data("out_lo", get_mem(get_single_element_layout(p), -127)),
+        data("out_hi", get_mem(get_single_element_layout(p), 127)),
+        softmax("softmax", "input", p.dimension),
+        reshape("reshape", "softmax", get_reshape_shape(p)),
+        quantize("quantize", "reshape", "in_lo", "in_hi", "out_lo", "out_hi", 255, data_types::i8),
+        reorder("reorder_bfyx", "quantize", get_output_layout(p))
+    );
+
+    tolerance = 1.f;
+    execute(p, {{"softmax", {"quantize"}}});
+}
+
+TEST_P(softmax_quantize_fusing_through, reorder) {
+    auto p = GetParam();
+    auto reorder_layout = layout{ p.data_type, p.input_format, get_reshape_shape(p), padding{} };
+    create_topologies(input_layout("input", get_input_layout(p)),
+        data("in_lo", get_mem(get_single_element_layout(p), min_random, 0)),
+        data("in_hi", get_mem(get_single_element_layout(p), 1, max_random)),
+        data("out_lo", get_mem(get_single_element_layout(p), -127)),
+        data("out_hi", get_mem(get_single_element_layout(p), 127)),
+        softmax("softmax", "input", p.dimension),
+        reorder("reorder", "softmax", reorder_layout),
+        quantize("quantize", "reorder", "in_lo", "in_hi", "out_lo", "out_hi", 255, data_types::i8),
+        reorder("reorder_bfyx", "quantize", get_output_layout(p))
+    );
+
+    tolerance = 1.f;
+    execute(p, {{"softmax", {"quantize"}}});
+}
+
+TEST_P(softmax_quantize_fusing_through, chain) {
+    auto p = GetParam();
+    auto reorder_layout = layout{ p.data_type, p.input_format, get_reshape_shape(p), padding{} };
+    create_topologies(input_layout("input", get_input_layout(p)),
+        data("in_lo", get_mem(get_single_element_layout(p), min_random, 0)),
+        data("in_hi", get_mem(get_single_element_layout(p), 1, max_random)),
+        data("out_lo", get_mem(get_single_element_layout(p), -127)),
+        data("out_hi", get_mem(get_single_element_layout(p), 127)),
+        softmax("softmax", "input", p.dimension),
+        reshape("reshape_first", "softmax", get_reshape_shape(p)),
+        reorder("reorder", "reshape_first", reorder_layout),
+        reshape("reshape_second", "reorder", get_reshape_shape(p)),
+        quantize("quantize", "reshape_second", "in_lo", "in_hi", "out_lo", "out_hi", 255, data_types::i8),
+        reorder("reorder_bfyx", "quantize", get_output_layout(p))
+    );
+
+    tolerance = 1.f;
+    execute(p, {{"softmax", {"quantize"}}});
+}
+
+INSTANTIATE_TEST_SUITE_P(fusings_gpu, softmax_quantize_fusing_through,
+    ::testing::ValuesIn(std::vector<softmax_test_params>{
+                        softmax_test_params{ CASE_SOFTMAX_FP32_1, 2, 3 },
+                        // Such fusions not allowed yet
+                        // softmax_test_params{ CASE_SOFTMAX_FP32_2, 3, 3 },
+                        // softmax_test_params{ CASE_SOFTMAX_FP32_3, 3, 3 },
+
+                        softmax_test_params{ CASE_SOFTMAX_FP16_1, 2, 3 },
+                        // Such fusions not allowed yet
+                        // softmax_test_params{ CASE_SOFTMAX_FP16_2, 3, 3 },
+                        // softmax_test_params{ CASE_SOFTMAX_FP16_3, 3, 3 },
+}));