[GPU] Add int8 fsv32 opt concat kernel (#34624)

### Description of the issue(symptom, root-cause, how it was resolved)
#### Symptom
The YOLO26n INT8 model shows a significant performance bottleneck on
GPU. Profiling revealed that feature-axis concatenation layers in
[b_fs_yx_fsv32](vscode-file://vscode-app/c:/Users/hyunback/AppData/Local/Programs/Microsoft%20VS%20Code/c3a26841a8/resources/app/out/vs/code/electron-browser/workbench/workbench.html)
format were falling back to the reference kernel, causing high latency.

#### Root cause
There was no optimized concatenation kernel for
[b_fs_yx_fsv32](vscode-file://vscode-app/c:/Users/hyunback/AppData/Local/Programs/Microsoft%20VS%20Code/c3a26841a8/resources/app/out/vs/code/electron-browser/workbench/workbench.html)
format. When the DPAS pipeline produces INT8 tensors in this format with
feature counts not aligned to 32, every concat operation dispatches the
generic reference kernel, which processes elements one by one without
leveraging sub-group block I/O.

#### Resolution
Added a new optimized concatenation kernel
(concatenation_gpu_b_fs_yx_fsv32) for feature-axis concat in
[b_fs_yx_fsv32](vscode-file://vscode-app/c:/Users/hyunback/AppData/Local/Programs/Microsoft%20VS%20Code/c3a26841a8/resources/app/out/vs/code/electron-browser/workbench/workbench.html)
format, supporting INT8/UINT8/FP16 data types.

The kernel handles two cases based on the cumulative output offset
alignment:

Aligned (output offset % 32 == 0): Uses block_read2/block_write2 for
maximum throughput. A runtime guard handles the last partial feature
block with scalar writes.
Unaligned (output offset % 32 != 0): Uses block_read2 for input +
per-element scalar writes with OUTPUT_GET_INDEX to correctly address
features that span across different FSV32 slices in memory.

#### Reproduction step and snapshot (if applicable. Do not attach for
customer model)
$ ./benchmark_app -m ./yolo26n_int8_openvino_model/yolo26n.xml -d GPU -t
10

#### Checklist 
 - [x] Is it a proper fix?
 - [x] Did you include test case for this fix, if necessary? 
- [x] Did you review existing test that can be extended to cover this
scenario? Which test did you review?

### Tickets:
 - *CVS-182396*

Signed-off-by: hyunback <hyunback.kim@intel.com>

Author: hyunback

src/plugins/intel_gpu/src/kernel_selector/cl_kernels/concatenation_gpu_b_fs_yx_fsv32.cl

@@ -0,0 +1,68 @@
+// Copyright (C) 2018-2026 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include "include/batch_headers/fetch_data.cl"
+#include "include/batch_headers/sub_group_block_read.cl"
+#include "include/batch_headers/sub_group_block_write.cl"
+
+// ======================================================================================
+// Optimized concatenation kernel for b_fs_yx_fsv32 format (feature-axis concat)
+// Supports INT8/UINT8/FP16 data types.
+//
+// Required JIT definitions:
+// --------------------------------------------------------------------------------------
+// SUB_GROUP_SIZE     - [int] sub-group/simd size; limited to 16
+// FSV                - [int] feature slice size; limited to 32
+// FSV_PER_THREAD     - [int] number of features per thread = FSV / SUB_GROUP_SIZE
+// ALIGNED            - [0/1] whether the output offset is aligned to FSV
+// ======================================================================================
+
+REQD_SUB_GROUP_SIZE(SUB_GROUP_SIZE)
+__attribute__((reqd_work_group_size(1, 1, SUB_GROUP_SIZE)))
+KERNEL(concatenation_gpu_b_fs_yx_fsv32)(
+    __global INPUT0_TYPE* input,
+    __global OUTPUT_TYPE* output,
+    uint output_offset_in_concat_axis)
+{
+    const uint x = (uint)get_global_id(0);
+    const uint y = (uint)get_global_id(1);
+    const uint fs_b_id = get_group_id(2);
+    const uint sglid = get_sub_group_local_id();
+
+    const uint fs = fs_b_id / INPUT0_BATCH_NUM;
+    const uint b = fs_b_id - fs * INPUT0_BATCH_NUM;
+
+    const uint input_offset = INPUT0_GET_INDEX(b, fs * FSV, y, x);
+
+    MAKE_VECTOR_TYPE(INPUT0_TYPE, 2) in = DT_INPUT_BLOCK_READ2(input, input_offset);
+
+    in = ACTIVATION(in, ACTIVATION_PARAMS);
+
+#if ALIGNED
+    const uint dst_index = OUTPUT_GET_INDEX(b, output_offset_in_concat_axis + fs * FSV, y, x);
+
+    // Full feature block: use block write for maximum throughput
+    if (fs * FSV + FSV <= INPUT0_FEATURE_NUM) {
+        DT_OUTPUT_BLOCK_WRITE2(output, dst_index, in);
+    } else {
+        // Last partial feature block: write only valid features
+        if (sglid + fs * FSV < INPUT0_FEATURE_NUM) {
+            output[dst_index + sglid] = in.s0;
+        }
+        if (sglid + SUB_GROUP_SIZE + fs * FSV < INPUT0_FEATURE_NUM) {
+            output[dst_index + SUB_GROUP_SIZE + sglid] = in.s1;
+        }
+    }
+#else
+    // Unaligned case: use per-element writes with proper index computation
+    const uint dst_feature = fs * FSV + output_offset_in_concat_axis + sglid;
+
+    if (sglid + SUB_GROUP_SIZE + fs * FSV < INPUT0_FEATURE_NUM) {
+        output[OUTPUT_GET_INDEX(b, dst_feature, y, x)] = in.s0;
+        output[OUTPUT_GET_INDEX(b, dst_feature + SUB_GROUP_SIZE, y, x)] = in.s1;
+    } else if (sglid + fs * FSV < INPUT0_FEATURE_NUM) {
+        output[OUTPUT_GET_INDEX(b, dst_feature, y, x)] = in.s0;
+    }
+#endif
+}

src/plugins/intel_gpu/src/kernel_selector/kernels/concatenation/concatenation_kernel_b_fs_yx_fsv32.cpp

@@ -0,0 +1,140 @@
+// Copyright (C) 2018-2026 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include <algorithm>
+#include "concatenation_kernel_b_fs_yx_fsv32.h"
+#include "kernel_selector_utils.h"
+
+namespace kernel_selector {
+
+static constexpr size_t subGroupSize = 16;
+static constexpr size_t fsv = 32;
+static constexpr size_t fsvPerThread = fsv / subGroupSize;
+
+ParamsKey ConcatenationKernel_b_fs_yx_fsv32::GetSupportedKey() const {
+    ParamsKey k;
+    k.EnableInputDataType(Datatype::INT8);
+    k.EnableOutputDataType(Datatype::INT8);
+    k.EnableInputDataType(Datatype::UINT8);
+    k.EnableOutputDataType(Datatype::UINT8);
+    k.EnableInputDataType(Datatype::F16);
+    k.EnableOutputDataType(Datatype::F16);
+    k.EnableInputLayout(DataLayout::b_fs_yx_fsv32);
+    k.EnableOutputLayout(DataLayout::b_fs_yx_fsv32);
+    k.EnableTensorOffset();
+    k.EnableTensorPitches();
+    k.EnableBatching();
+    k.EnableConcatAxis(ConcatAxis::FEATURE);
+    k.EnableConcatKernelPerInput();
+    return k;
+}
+
+DeviceFeaturesKey ConcatenationKernel_b_fs_yx_fsv32::get_required_device_features_key(const Params& params) const {
+    return get_common_subgroups_device_features_key(params);
+}
+
+bool ConcatenationKernel_b_fs_yx_fsv32::Validate(const Params& p) const {
+    if (!ConcatenationKernelBase::Validate(p)) {
+        DO_NOT_USE_THIS_KERNEL(p.layerID);
+    }
+
+    const concatenation_params& params = static_cast<const concatenation_params&>(p);
+
+    if (params.axis != ConcatAxis::FEATURE)
+        DO_NOT_USE_THIS_KERNEL(p.layerID);
+
+    // All inputs must have the same layout
+    auto same_layout = params.inputs[0].GetLayout();
+    for (const auto& lt : params.inputs) {
+        if (lt.GetLayout() != same_layout) {
+            DO_NOT_USE_THIS_KERNEL(p.layerID);
+        }
+    }
+
+    return true;
+}
+
+ConcatenationKernelBase::DispatchData ConcatenationKernel_b_fs_yx_fsv32::SetDefault(const concatenation_params& params) const {
+    DispatchData dispatchData = ConcatenationKernelBase::SetDefault(params);
+    const auto& input = params.inputs[0];
+
+    dispatchData.gws[0] = input.X().v;
+    dispatchData.gws[1] = input.Y().v;
+    dispatchData.gws[2] = CeilDiv(input.Feature().v, fsv) * subGroupSize * input.Batch().v;
+
+    dispatchData.lws[0] = 1;
+    dispatchData.lws[1] = 1;
+    dispatchData.lws[2] = subGroupSize;
+
+    return dispatchData;
+}
+
+KernelsPriority ConcatenationKernel_b_fs_yx_fsv32::GetKernelsPriority(const Params& /*params*/) const {
+    return FORCE_PRIORITY_1;
+}
+
+JitConstants ConcatenationKernel_b_fs_yx_fsv32::GetJitConstants(const concatenation_params& params) const {
+    JitConstants jit = MakeBaseParamsJitConstants(params);
+
+    jit.AddConstant(MakeJitConstant("ALIGNED", params.isAligned));
+    jit.AddConstant(MakeJitConstant("FSV", fsv));
+    jit.AddConstant(MakeJitConstant("SUB_GROUP_SIZE", subGroupSize));
+    jit.AddConstant(MakeJitConstant("FSV_PER_THREAD", fsvPerThread));
+
+    return jit;
+}
+
+size_t ConcatenationKernel_b_fs_yx_fsv32::GetAlignment(const concatenation_params& /*params*/) const {
+    return fsv;
+}
+
+KernelsData ConcatenationKernel_b_fs_yx_fsv32::GetKernelsData(const Params& params) const {
+    if (!Validate(params)) {
+        return {};
+    }
+
+    const concatenation_params& orgParams = static_cast<const concatenation_params&>(params);
+
+    KernelData kd = KernelData::Default<concatenation_params>(params, orgParams.inputs.size());
+
+    uint32_t lastOffset = 0;
+    size_t ifm_offset = 0;
+    for (size_t i = 0; i < orgParams.inputs.size(); i++) {
+        const auto& input = orgParams.inputs[i];
+
+        auto newParams = orgParams;
+        newParams.inputs.resize(1);
+        newParams.inputs[0] = input;
+        size_t ifm = input.Feature().v;
+        newParams.isAligned = ifm_offset % fsv == 0;
+        ifm_offset += ifm;
+
+        auto& kernel = kd.kernels[i];
+        DispatchData dispatchData = SetDefault(newParams);
+        auto cldnnJit = GetJitConstants(newParams);
+        auto entryPoint = GetEntryPoint(kernelName, newParams.layerID, params, i);
+        auto jit = CreateJit(kernelName, cldnnJit, entryPoint);
+
+        kernel.code.kernelString = GetKernelString(kernelName, jit, entryPoint, params.engineInfo);
+        kernel.params.workGroups.global = dispatchData.gws;
+        kernel.params.workGroups.local = dispatchData.lws;
+        kernel.params.arguments.push_back({ArgumentDescriptor::Types::INPUT, (uint32_t)i});
+        kernel.params.arguments.push_back({ArgumentDescriptor::Types::OUTPUT, 0});
+        kernel.skip_execution = KernelData::SkipKernelExecution(newParams);
+
+        ScalarDescriptor s;
+        s.t = ScalarDescriptor::Types::UINT32;
+        s.v.u32 = lastOffset;
+        kernel.params.scalars.push_back(s);
+        kernel.params.arguments.push_back({ArgumentDescriptor::Types::SCALAR, 0});
+
+        auto concatChannelIndex = DataTensor::Channelndex(orgParams.inputs[i].GetLayout(), GetConcatChannel(orgParams));
+        OPENVINO_ASSERT(concatChannelIndex >= 0, "concatChannelIndex shouldn't be negative");
+        lastOffset += (uint32_t)input.GetDims()[concatChannelIndex].v;
+    }
+
+    return {kd};
+}
+
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/kernels/concatenation/concatenation_kernel_b_fs_yx_fsv32.h

@@ -0,0 +1,26 @@
+// Copyright (C) 2018-2026 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma once
+
+#include "concatenation_kernel_base.h"
+
+namespace kernel_selector {
+
+class ConcatenationKernel_b_fs_yx_fsv32 : public ConcatenationKernelBase {
+public:
+    ConcatenationKernel_b_fs_yx_fsv32() : ConcatenationKernelBase("concatenation_gpu_b_fs_yx_fsv32") {}
+    virtual ~ConcatenationKernel_b_fs_yx_fsv32() {}
+
+    KernelsData GetKernelsData(const Params& params) const override;
+    KernelsPriority GetKernelsPriority(const Params& params) const override;
+    ParamsKey GetSupportedKey() const override;
+    DeviceFeaturesKey get_required_device_features_key(const Params& params) const override;
+    DispatchData SetDefault(const concatenation_params& params) const override;
+    JitConstants GetJitConstants(const concatenation_params& params) const override;
+    bool Validate(const Params& p) const override;
+    size_t GetAlignment(const concatenation_params& params) const override;
+};
+
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/kernels/concatenation/concatenation_kernel_selector.cpp

@@ -7,6 +7,7 @@
 #include "concatenation_kernel_simple_ref.h"
 #include "concatenation_kernel_depth_bfyx_no_pitch.h"
 #include "concatenation_kernel_b_fs_yx_fsv16.h"
+#include "concatenation_kernel_b_fs_yx_fsv32.h"
 #include "concatenation_kernel_fs_b_yx_fsv32.h"
 
 namespace kernel_selector {
@@ -15,6 +16,7 @@ concatenation_kernel_selector::concatenation_kernel_selector() {
     Attach<ConcatenationKernel_simple_Ref>();
     Attach<ConcatenationKernel_depth_bfyx_no_pitch>();
     Attach<ConcatenationKernel_b_fs_yx_fsv16>();
+    Attach<ConcatenationKernel_b_fs_yx_fsv32>();
     Attach<ConcatenationKernel_fs_b_yx_fsv32>();
 }
 

src/plugins/intel_gpu/tests/unit/test_cases/concatenation_gpu_test.cpp

@@ -2182,3 +2182,135 @@ INSTANTIATE_TEST_SUITE_P(smoke,
                         ),
                         concat_gpu_implicit::PrintToStringParamName);
 #endif
+// ============================================================================
+// Tests for concatenation_gpu_b_fs_yx_fsv32 kernel (int8/uint8, feature axis)
+// Forces OCL impl with b_fs_yx_fsv32 format to ensure the optimized kernel
+// is selected instead of falling back to ref kernel.
+// ============================================================================
+using TestParamType_concat_fsv32 = ::testing::tuple<
+    size_t,                   // 0 - Batch
+    std::vector<size_t>,      // 1 - Input feature sizes
+    size_t,                   // 2 - Input Y
+    size_t,                   // 3 - Input X
+    data_types>;              // 4 - Data type (i8/u8)
+
+struct concat_gpu_b_fs_yx_fsv32_force : public ::testing::TestWithParam<TestParamType_concat_fsv32> {
+    tests::random_generator rg;
+
+    void SetUp() override {
+        rg.set_seed(GET_SUITE_NAME);
+    }
+
+    static std::string PrintToStringParamName(testing::TestParamInfo<TestParamType_concat_fsv32> param_info) {
+        auto batch = testing::get<0>(param_info.param);
+        auto& feats = testing::get<1>(param_info.param);
+        auto y = testing::get<2>(param_info.param);
+        auto x = testing::get<3>(param_info.param);
+        auto dt = testing::get<4>(param_info.param);
+        std::string feat_str;
+        for (size_t i = 0; i < feats.size(); i++)
+            feat_str += (i ? "_" : "") + std::to_string(feats[i]);
+        return "b" + std::to_string(batch) + "_f" + feat_str +
+               "_y" + std::to_string(y) + "_x" + std::to_string(x) +
+               "_" + ov::element::Type(dt).get_type_name();
+    }
+
+    void test() {
+        auto& engine = get_test_engine();
+        const auto batch_num = testing::get<0>(GetParam());
+        const auto& in_features = testing::get<1>(GetParam());
+        const auto input_y = testing::get<2>(GetParam());
+        const auto input_x = testing::get<3>(GetParam());
+        const auto dt = testing::get<4>(GetParam());
+
+        topology topology;
+        std::vector<memory::ptr> in_memory;
+        std::vector<input_info> input_ids;
+
+        for (size_t i = 0; i < in_features.size(); i++) {
+            auto sz = tensor(static_cast<int32_t>(batch_num), static_cast<int32_t>(in_features[i]),
+                             static_cast<int32_t>(input_x), static_cast<int32_t>(input_y));
+            auto in_lay = layout(dt, format::b_fs_yx_fsv32, sz);
+            auto in_mem = engine.allocate_memory(in_lay);
+            auto count = static_cast<int>(batch_num * in_features[i] * input_y * input_x);
+            if (dt == data_types::i8)
+                set_values<int8_t>(in_mem, rg.generate_random_1d<int8_t>(count, -50, 50));
+            else
+                set_values<uint8_t>(in_mem, rg.generate_random_1d<uint8_t>(count, 0, 200));
+            in_memory.push_back(in_mem);
+            topology.add(input_layout("input" + std::to_string(i), in_lay));
+            input_ids.push_back(input_info("input" + std::to_string(i)));
+        }
+
+        topology.add(concatenation("concat", input_ids, 1, dt));
+        topology.add(reorder("output", input_info("concat"), format::bfyx, dt));
+
+        ExecutionConfig config = get_test_default_config(engine);
+        config.set_property(ov::intel_gpu::optimize_data(true));
+        ov::intel_gpu::ImplementationDesc impl = { format::b_fs_yx_fsv32, "concatenation_gpu_b_fs_yx_fsv32", impl_types::ocl };
+        config.set_property(ov::intel_gpu::force_implementations(ov::intel_gpu::ImplForcingMap{ { "concat", impl } }));
+
+        network network(engine, topology, config);
+        for (size_t i = 0; i < in_features.size(); i++)
+            network.set_input_data(input_ids[i].pid, in_memory[i]);
+
+        auto outputs = network.execute();
+        auto out_mem = outputs.at("output").get_memory();
+        auto out_lay = out_mem->get_layout();
+        size_t total_f = 0;
+        for (auto f : in_features) total_f += f;
+        ASSERT_EQ(out_lay.feature(), static_cast<int>(total_f));
+
+        // Verify element-by-element against inputs read in logical order
+        size_t f_offset = 0;
+        for (size_t in_i = 0; in_i < in_features.size(); in_i++) {
+            for (size_t b = 0; b < batch_num; b++) {
+                for (size_t f = 0; f < in_features[in_i]; f++) {
+                    for (size_t y = 0; y < input_y; y++) {
+                        for (size_t x = 0; x < input_x; x++) {
+                            auto in_coords = tensor(batch(b), feature(f), spatial(x, y, 0, 0));
+                            auto out_coords = tensor(batch(b), feature(f_offset + f), spatial(x, y, 0, 0));
+                            if (dt == data_types::i8) {
+                                cldnn::mem_lock<int8_t> in_ptr(in_memory[in_i], get_test_stream());
+                                cldnn::mem_lock<int8_t> out_ptr(out_mem, get_test_stream());
+                                ASSERT_EQ(in_ptr[in_memory[in_i]->get_layout().get_linear_offset(in_coords)],
+                                          out_ptr[out_lay.get_linear_offset(out_coords)])
+                                    << " b=" << b << " f=" << f_offset + f << " y=" << y << " x=" << x;
+                            } else {
+                                cldnn::mem_lock<uint8_t> in_ptr(in_memory[in_i], get_test_stream());
+                                cldnn::mem_lock<uint8_t> out_ptr(out_mem, get_test_stream());
+                                ASSERT_EQ(in_ptr[in_memory[in_i]->get_layout().get_linear_offset(in_coords)],
+                                          out_ptr[out_lay.get_linear_offset(out_coords)])
+                                    << " b=" << b << " f=" << f_offset + f << " y=" << y << " x=" << x;
+                            }
+                        }
+                    }
+                }
+            }
+            f_offset += in_features[in_i];
+        }
+    }
+};
+
+TEST_P(concat_gpu_b_fs_yx_fsv32_force, feature_axis) {
+    ASSERT_NO_FATAL_FAILURE(test());
+}
+
+INSTANTIATE_TEST_SUITE_P(smoke,
+    concat_gpu_b_fs_yx_fsv32_force,
+    ::testing::Values(
+        // Aligned features (all 32-aligned)
+        TestParamType_concat_fsv32(1, { 32, 64, 32 }, 2, 2, data_types::i8),
+        TestParamType_concat_fsv32(1, { 32, 64, 32 }, 2, 2, data_types::u8),
+        // Unaligned features
+        TestParamType_concat_fsv32(1, { 24, 48, 17 }, 3, 3, data_types::i8),
+        TestParamType_concat_fsv32(1, { 24, 48, 17 }, 3, 3, data_types::u8),
+        // Mixed aligned/unaligned, batch > 1
+        TestParamType_concat_fsv32(2, { 64, 33, 15 }, 2, 2, data_types::i8),
+        TestParamType_concat_fsv32(2, { 64, 33, 15 }, 2, 2, data_types::u8),
+        // Single small unaligned input
+        TestParamType_concat_fsv32(1, { 3, 5 }, 4, 4, data_types::i8),
+        // Large aligned
+        TestParamType_concat_fsv32(1, { 64, 64, 64, 64 }, 1, 1, data_types::i8)
+    ),
+    concat_gpu_b_fs_yx_fsv32_force::PrintToStringParamName);