diff --git a/dpnp/backend/extensions/indexing/CMakeLists.txt b/dpnp/backend/extensions/indexing/CMakeLists.txt
index 370d59f95585..e1bc34c9ae8b 100644
--- a/dpnp/backend/extensions/indexing/CMakeLists.txt
+++ b/dpnp/backend/extensions/indexing/CMakeLists.txt
@@ -62,7 +62,7 @@ set_target_properties(
 
 target_include_directories(
     ${python_module_name}
-    PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../common
+    PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../ ${CMAKE_CURRENT_SOURCE_DIR}/../common
 )
 
 # treat below headers as system to suppress the warnings there during the build
diff --git a/dpnp/backend/extensions/indexing/choose.cpp b/dpnp/backend/extensions/indexing/choose.cpp
index 99d91744366f..91dcf6b00213 100644
--- a/dpnp/backend/extensions/indexing/choose.cpp
+++ b/dpnp/backend/extensions/indexing/choose.cpp
@@ -30,41 +30,116 @@
 #include <cstddef>
 #include <cstdint>
 #include <memory>
-#include <pybind11/pybind11.h>
-#include <pybind11/stl.h>
-#include <sycl/sycl.hpp>
 #include <type_traits>
 #include <utility>
 #include <vector>
 
-#include "choose_kernel.hpp"
+#include <sycl/sycl.hpp>
+
 #include "dpctl4pybind11.hpp"
+#include <pybind11/pybind11.h>
+#include <pybind11/stl.h>
 
-// utils extension header
 #include "ext/common.hpp"
+#include "kernels/indexing/choose.hpp"
 
 // dpctl tensor headers
 #include "utils/indexing_utils.hpp"
 #include "utils/memory_overlap.hpp"
+#include "utils/offset_utils.hpp" //
 #include "utils/output_validation.hpp"
 #include "utils/sycl_alloc_utils.hpp"
 #include "utils/type_dispatch.hpp"
 
 namespace dpnp::extensions::indexing
 {
-
+namespace py = pybind11;
 namespace td_ns = dpctl::tensor::type_dispatch;
 
-static kernels::choose_fn_ptr_t choose_clip_dispatch_table[td_ns::num_types]
-                                                          [td_ns::num_types];
-static kernels::choose_fn_ptr_t choose_wrap_dispatch_table[td_ns::num_types]
-                                                          [td_ns::num_types];
+using dpctl::tensor::ssize_t;
+
+typedef sycl::event (*choose_fn_ptr_t)(sycl::queue &,
+                                       size_t,
+                                       ssize_t,
+                                       int,
+                                       const ssize_t *,
+                                       const char *,
+                                       char *,
+                                       char **,
+                                       ssize_t,
+                                       ssize_t,
+                                       const ssize_t *,
+                                       const std::vector<sycl::event> &);
+
+static choose_fn_ptr_t choose_clip_dispatch_table[td_ns::num_types]
+                                                 [td_ns::num_types];
+static choose_fn_ptr_t choose_wrap_dispatch_table[td_ns::num_types]
+                                                 [td_ns::num_types];
+
+template <typename ProjectorT, typename indTy, typename Ty>
+sycl::event choose_impl(sycl::queue &q,
+                        size_t nelems,
+                        ssize_t n_chcs,
+                        int nd,
+                        const ssize_t *shape_and_strides,
+                        const char *ind_cp,
+                        char *dst_cp,
+                        char **chcs_cp,
+                        ssize_t ind_offset,
+                        ssize_t dst_offset,
+                        const ssize_t *chc_offsets,
+                        const std::vector<sycl::event> &depends)
+{
+    dpctl::tensor::type_utils::validate_type_for_device<Ty>(q);
+
+    const indTy *ind_tp = reinterpret_cast<const indTy *>(ind_cp);
+    Ty *dst_tp = reinterpret_cast<Ty *>(dst_cp);
 
-namespace py = pybind11;
+    sycl::event choose_ev = q.submit([&](sycl::handler &cgh) {
+        cgh.depends_on(depends);
 
-namespace detail
+        using InOutIndexerT =
+            dpctl::tensor::offset_utils::TwoOffsets_StridedIndexer;
+        const InOutIndexerT ind_out_indexer{nd, ind_offset, dst_offset,
+                                            shape_and_strides};
+
+        using NthChoiceIndexerT =
+            dpnp::kernels::choose::strides::NthStrideOffsetUnpacked;
+        const NthChoiceIndexerT choices_indexer{
+            nd, chc_offsets, shape_and_strides, shape_and_strides + 3 * nd};
+
+        using ChooseFunc =
+            dpnp::kernels::choose::ChooseFunctor<ProjectorT, InOutIndexerT,
+                                                 NthChoiceIndexerT, indTy, Ty>;
+
+        cgh.parallel_for<ChooseFunc>(sycl::range<1>(nelems),
+                                     ChooseFunc(ind_tp, dst_tp, chcs_cp, n_chcs,
+                                                ind_out_indexer,
+                                                choices_indexer));
+    });
+
+    return choose_ev;
+}
+
+template <typename fnT, typename IndT, typename T, typename Index>
+struct ChooseFactory
 {
+    fnT get()
+    {
+        if constexpr (std::is_integral<IndT>::value &&
+                      !std::is_same<IndT, bool>::value) {
+            fnT fn = choose_impl<Index, IndT, T>;
+            return fn;
+        }
+        else {
+            fnT fn = nullptr;
+            return fn;
+        }
+    }
+};
 
+namespace detail
+{
 using host_ptrs_allocator_t =
     dpctl::tensor::alloc_utils::usm_host_allocator<char *>;
 using ptrs_t = std::vector<char *, host_ptrs_allocator_t>;
@@ -191,7 +266,6 @@ std::vector<dpctl::tensor::usm_ndarray> parse_py_chcs(const sycl::queue &q,
 
     return res;
 }
-
 } // namespace detail
 
 std::pair<sycl::event, sycl::event>
@@ -412,23 +486,6 @@ std::pair<sycl::event, sycl::event>
     return std::make_pair(arg_cleanup_ev, choose_generic_ev);
 }
 
-template <typename fnT, typename IndT, typename T, typename Index>
-struct ChooseFactory
-{
-    fnT get()
-    {
-        if constexpr (std::is_integral<IndT>::value &&
-                      !std::is_same<IndT, bool>::value) {
-            fnT fn = kernels::choose_impl<Index, IndT, T>;
-            return fn;
-        }
-        else {
-            fnT fn = nullptr;
-            return fn;
-        }
-    }
-};
-
 using dpctl::tensor::indexing_utils::ClipIndex;
 using dpctl::tensor::indexing_utils::WrapIndex;
 
@@ -441,7 +498,6 @@ using ChooseClipFactory = ChooseFactory<fnT, IndT, T, ClipIndex<IndT>>;
 void init_choose_dispatch_tables(void)
 {
     using ext::common::init_dispatch_table;
-    using kernels::choose_fn_ptr_t;
 
     init_dispatch_table<choose_fn_ptr_t, ChooseClipFactory>(
         choose_clip_dispatch_table);
diff --git a/dpnp/backend/extensions/indexing/choose_kernel.hpp b/dpnp/backend/extensions/indexing/choose_kernel.hpp
deleted file mode 100644
index 6b1ac8005054..000000000000
--- a/dpnp/backend/extensions/indexing/choose_kernel.hpp
+++ /dev/null
@@ -1,191 +0,0 @@
-//*****************************************************************************
-// Copyright (c) 2025, Intel Corporation
-// All rights reserved.
-//
-// Redistribution and use in source and binary forms, with or without
-// modification, are permitted provided that the following conditions are met:
-// - Redistributions of source code must retain the above copyright notice,
-//   this list of conditions and the following disclaimer.
-// - Redistributions in binary form must reproduce the above copyright notice,
-//   this list of conditions and the following disclaimer in the documentation
-//   and/or other materials provided with the distribution.
-// - Neither the name of the copyright holder nor the names of its contributors
-//   may be used to endorse or promote products derived from this software
-//   without specific prior written permission.
-//
-// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
-// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
-// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
-// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
-// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
-// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
-// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
-// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
-// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
-// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
-// THE POSSIBILITY OF SUCH DAMAGE.
-//*****************************************************************************
-
-#pragma once
-
-#include <algorithm>
-#include <complex>
-#include <cstdint>
-#include <limits>
-#include <type_traits>
-
-#include <sycl/sycl.hpp>
-
-#include "kernels/dpctl_tensor_types.hpp"
-#include "utils/indexing_utils.hpp"
-#include "utils/offset_utils.hpp"
-#include "utils/strided_iters.hpp"
-#include "utils/type_utils.hpp"
-
-namespace dpnp::extensions::indexing::strides_detail
-{
-
-struct NthStrideOffsetUnpacked
-{
-    NthStrideOffsetUnpacked(int common_nd,
-                            dpctl::tensor::ssize_t const *_offsets,
-                            dpctl::tensor::ssize_t const *_shape,
-                            dpctl::tensor::ssize_t const *_strides)
-        : _ind(common_nd), nd(common_nd), offsets(_offsets), shape(_shape),
-          strides(_strides)
-    {
-    }
-
-    template <typename nT>
-    size_t operator()(dpctl::tensor::ssize_t gid, nT n) const
-    {
-        dpctl::tensor::ssize_t relative_offset(0);
-        _ind.get_displacement<const dpctl::tensor::ssize_t *,
-                              const dpctl::tensor::ssize_t *>(
-            gid, shape, strides + (n * nd), relative_offset);
-
-        return relative_offset + offsets[n];
-    }
-
-private:
-    dpctl::tensor::strides::CIndexer_vector<dpctl::tensor::ssize_t> _ind;
-
-    int nd;
-    dpctl::tensor::ssize_t const *offsets;
-    dpctl::tensor::ssize_t const *shape;
-    dpctl::tensor::ssize_t const *strides;
-};
-
-static_assert(sycl::is_device_copyable_v<NthStrideOffsetUnpacked>);
-
-} // namespace dpnp::extensions::indexing::strides_detail
-
-namespace dpnp::extensions::indexing::kernels
-{
-
-template <typename ProjectorT,
-          typename IndOutIndexerT,
-          typename ChoicesIndexerT,
-          typename IndT,
-          typename T>
-class ChooseFunctor
-{
-private:
-    const IndT *ind = nullptr;
-    T *dst = nullptr;
-    char **chcs = nullptr;
-    dpctl::tensor::ssize_t n_chcs;
-    const IndOutIndexerT ind_out_indexer;
-    const ChoicesIndexerT chcs_indexer;
-
-public:
-    ChooseFunctor(const IndT *ind_,
-                  T *dst_,
-                  char **chcs_,
-                  dpctl::tensor::ssize_t n_chcs_,
-                  const IndOutIndexerT &ind_out_indexer_,
-                  const ChoicesIndexerT &chcs_indexer_)
-        : ind(ind_), dst(dst_), chcs(chcs_), n_chcs(n_chcs_),
-          ind_out_indexer(ind_out_indexer_), chcs_indexer(chcs_indexer_)
-    {
-    }
-
-    void operator()(sycl::id<1> id) const
-    {
-        const ProjectorT proj{};
-
-        dpctl::tensor::ssize_t i = id[0];
-
-        auto ind_dst_offsets = ind_out_indexer(i);
-        dpctl::tensor::ssize_t ind_offset = ind_dst_offsets.get_first_offset();
-        dpctl::tensor::ssize_t dst_offset = ind_dst_offsets.get_second_offset();
-
-        IndT chc_idx = ind[ind_offset];
-        // proj produces an index in the range of n_chcs
-        dpctl::tensor::ssize_t projected_idx = proj(n_chcs, chc_idx);
-
-        dpctl::tensor::ssize_t chc_offset = chcs_indexer(i, projected_idx);
-
-        T *chc = reinterpret_cast<T *>(chcs[projected_idx]);
-
-        dst[dst_offset] = chc[chc_offset];
-    }
-};
-
-typedef sycl::event (*choose_fn_ptr_t)(sycl::queue &,
-                                       size_t,
-                                       dpctl::tensor::ssize_t,
-                                       int,
-                                       const dpctl::tensor::ssize_t *,
-                                       const char *,
-                                       char *,
-                                       char **,
-                                       dpctl::tensor::ssize_t,
-                                       dpctl::tensor::ssize_t,
-                                       const dpctl::tensor::ssize_t *,
-                                       const std::vector<sycl::event> &);
-
-template <typename ProjectorT, typename indTy, typename Ty>
-sycl::event choose_impl(sycl::queue &q,
-                        size_t nelems,
-                        dpctl::tensor::ssize_t n_chcs,
-                        int nd,
-                        const dpctl::tensor::ssize_t *shape_and_strides,
-                        const char *ind_cp,
-                        char *dst_cp,
-                        char **chcs_cp,
-                        dpctl::tensor::ssize_t ind_offset,
-                        dpctl::tensor::ssize_t dst_offset,
-                        const dpctl::tensor::ssize_t *chc_offsets,
-                        const std::vector<sycl::event> &depends)
-{
-    dpctl::tensor::type_utils::validate_type_for_device<Ty>(q);
-
-    const indTy *ind_tp = reinterpret_cast<const indTy *>(ind_cp);
-    Ty *dst_tp = reinterpret_cast<Ty *>(dst_cp);
-
-    sycl::event choose_ev = q.submit([&](sycl::handler &cgh) {
-        cgh.depends_on(depends);
-
-        using InOutIndexerT =
-            dpctl::tensor::offset_utils::TwoOffsets_StridedIndexer;
-        const InOutIndexerT ind_out_indexer{nd, ind_offset, dst_offset,
-                                            shape_and_strides};
-
-        using NthChoiceIndexerT = strides_detail::NthStrideOffsetUnpacked;
-        const NthChoiceIndexerT choices_indexer{
-            nd, chc_offsets, shape_and_strides, shape_and_strides + 3 * nd};
-
-        using ChooseFunc = ChooseFunctor<ProjectorT, InOutIndexerT,
-                                         NthChoiceIndexerT, indTy, Ty>;
-
-        cgh.parallel_for<ChooseFunc>(sycl::range<1>(nelems),
-                                     ChooseFunc(ind_tp, dst_tp, chcs_cp, n_chcs,
-                                                ind_out_indexer,
-                                                choices_indexer));
-    });
-
-    return choose_ev;
-}
-
-} // namespace dpnp::extensions::indexing::kernels
diff --git a/dpnp/backend/kernels/indexing/choose.hpp b/dpnp/backend/kernels/indexing/choose.hpp
new file mode 100644
index 000000000000..7129ca7bce1e
--- /dev/null
+++ b/dpnp/backend/kernels/indexing/choose.hpp
@@ -0,0 +1,126 @@
+//*****************************************************************************
+// Copyright (c) 2024, Intel Corporation
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+// - Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// - Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// - Neither the name of the copyright holder nor the names of its contributors
+//   may be used to endorse or promote products derived from this software
+//   without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
+// THE POSSIBILITY OF SUCH DAMAGE.
+//*****************************************************************************
+
+#pragma once
+
+#include <sycl/sycl.hpp>
+
+#include "kernels/dpctl_tensor_types.hpp"
+#include "utils/strided_iters.hpp"
+
+namespace dpnp::kernels::choose
+{
+using dpctl::tensor::ssize_t;
+
+template <typename ProjectorT,
+          typename IndOutIndexerT,
+          typename ChoicesIndexerT,
+          typename IndT,
+          typename T>
+class ChooseFunctor
+{
+private:
+    const IndT *ind = nullptr;
+    T *dst = nullptr;
+    char **chcs = nullptr;
+    ssize_t n_chcs;
+    const IndOutIndexerT ind_out_indexer;
+    const ChoicesIndexerT chcs_indexer;
+
+public:
+    ChooseFunctor(const IndT *ind_,
+                  T *dst_,
+                  char **chcs_,
+                  ssize_t n_chcs_,
+                  const IndOutIndexerT &ind_out_indexer_,
+                  const ChoicesIndexerT &chcs_indexer_)
+        : ind(ind_), dst(dst_), chcs(chcs_), n_chcs(n_chcs_),
+          ind_out_indexer(ind_out_indexer_), chcs_indexer(chcs_indexer_)
+    {
+    }
+
+    void operator()(sycl::id<1> id) const
+    {
+        const ProjectorT proj{};
+
+        ssize_t i = id[0];
+
+        auto ind_dst_offsets = ind_out_indexer(i);
+        ssize_t ind_offset = ind_dst_offsets.get_first_offset();
+        ssize_t dst_offset = ind_dst_offsets.get_second_offset();
+
+        IndT chc_idx = ind[ind_offset];
+        // proj produces an index in the range of n_chcs
+        ssize_t projected_idx = proj(n_chcs, chc_idx);
+
+        ssize_t chc_offset = chcs_indexer(i, projected_idx);
+
+        T *chc = reinterpret_cast<T *>(chcs[projected_idx]);
+
+        dst[dst_offset] = chc[chc_offset];
+    }
+};
+
+namespace strides
+{
+using dpctl::tensor::strides::CIndexer_vector;
+
+struct NthStrideOffsetUnpacked
+{
+    NthStrideOffsetUnpacked(int common_nd,
+                            ssize_t const *_offsets,
+                            ssize_t const *_shape,
+                            ssize_t const *_strides)
+        : _ind(common_nd), nd(common_nd), offsets(_offsets), shape(_shape),
+          strides(_strides)
+    {
+    }
+
+    template <typename nT>
+    size_t operator()(ssize_t gid, nT n) const
+    {
+        ssize_t relative_offset(0);
+        _ind.get_displacement<const ssize_t *, const ssize_t *>(
+            gid, shape, strides + (n * nd), relative_offset);
+
+        return relative_offset + offsets[n];
+    }
+
+private:
+    CIndexer_vector<ssize_t> _ind;
+
+    int nd;
+    ssize_t const *offsets;
+    ssize_t const *shape;
+    ssize_t const *strides;
+};
+
+static_assert(sycl::is_device_copyable_v<NthStrideOffsetUnpacked>);
+
+} // namespace strides
+} // namespace dpnp::kernels::choose
diff --git a/scripts/gen_coverage.py b/scripts/gen_coverage.py
index 588345d91b2e..d245c5d31060 100644
--- a/scripts/gen_coverage.py
+++ b/scripts/gen_coverage.py
@@ -259,6 +259,7 @@ def find_objects():
                     "-format=lcov",
                     "-ignore-filename-regex=/tmp/icpx*",
                     r"-ignore-filename-regex=.*/backend/kernels/elementwise_functions/.*\.hpp$",
+                    r"-ignore-filename-regex=.*/backend/kernels/indexing/.*\.hpp$",
                     "-instr-profile=" + instr_profile_fn,
                 ]
                 + objects