diff --git a/CHANGELOG.md b/CHANGELOG.md
index 6bc89e7b7bef..2556938f7797 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -36,6 +36,7 @@ This release is compatible with NumPy 2.5.
 * Replaced references to the `dpnp.amax`/`dpnp.amin` aliases with the canonical `dpnp.max`/`dpnp.min` in docstrings and code internally [#2990](https://github.com/IntelPython/dpnp/pull/2990)
 * Aligned the signature of `dpnp.tensor.expand_dims` with the Python array API by making `axis` a required argument [#2988](https://github.com/IntelPython/dpnp/pull/2988)
 * Removed dead code branches guarded by outdated oneMKL and DPC++ compiler version checks [#2999](https://github.com/IntelPython/dpnp/pull/2999)
+* Replaced the deprecated `nd_item::barrier()` member calls in the `accumulators` and `gemm` kernels with the SYCL 2020 `sycl::group_barrier()` free function [#3006](https://github.com/IntelPython/dpnp/pull/3006)
 
 ### Deprecated
 
diff --git a/dpnp/tensor/libtensor/include/kernels/accumulators.hpp b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp
index 2989cf26bb4e..079fcf5e9c96 100644
--- a/dpnp/tensor/libtensor/include/kernels/accumulators.hpp
+++ b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp
@@ -302,11 +302,12 @@ sycl::event inclusive_scan_base_step_blocked(
                     local_iscan.back(), identity, scan_op);
                 // ensure all finished reading from SLM, to avoid race condition
                 // with subsequent writes into SLM
-                it.barrier(sycl::access::fence_space::local_space);
+                sycl::group_barrier(it.get_group(),
+                                    sycl::memory_scope::work_group);
             }
 
             slm_iscan_tmp[(lid + 1) % wg_size] = wg_iscan_val;
-            it.barrier(sycl::access::fence_space::local_space);
+            sycl::group_barrier(it.get_group(), sycl::memory_scope::work_group);
             const outputT modifier = (lid == 0) ? identity : slm_iscan_tmp[lid];
 
 #pragma unroll
@@ -430,7 +431,8 @@ sycl::event inclusive_scan_base_step_striped(
                         slm_iscan_tmp[local_offset0 + i] = local_iscan[i];
                     }
 
-                    it.barrier(sycl::access::fence_space::local_space);
+                    sycl::group_barrier(it.get_group(),
+                                        sycl::memory_scope::work_group);
                 }
 
                 {
@@ -452,7 +454,8 @@ sycl::event inclusive_scan_base_step_striped(
                             slm_iscan_tmp[block_offset + disp_exchanged];
                     }
 
-                    it.barrier(sycl::access::fence_space::local_space);
+                    sycl::group_barrier(it.get_group(),
+                                        sycl::memory_scope::work_group);
                 }
             }
 
@@ -476,11 +479,12 @@ sycl::event inclusive_scan_base_step_striped(
                     identity, scan_op);
                 // ensure all finished reading from SLM, to avoid race condition
                 // with subsequent writes into SLM
-                it.barrier(sycl::access::fence_space::local_space);
+                sycl::group_barrier(it.get_group(),
+                                    sycl::memory_scope::work_group);
             }
 
             slm_iscan_tmp[(lid + 1) % wg_size] = wg_iscan_val;
-            it.barrier(sycl::access::fence_space::local_space);
+            sycl::group_barrier(it.get_group(), sycl::memory_scope::work_group);
             const outputT modifier = (lid == 0) ? identity : slm_iscan_tmp[lid];
 
 #pragma unroll
@@ -488,7 +492,7 @@ sycl::event inclusive_scan_base_step_striped(
                 local_iscan[m_wi] = scan_op(local_iscan[m_wi], modifier);
             }
 
-            it.barrier(sycl::access::fence_space::local_space);
+            sycl::group_barrier(it.get_group(), sycl::memory_scope::work_group);
 
             // convert back to blocked layout
             {
@@ -499,7 +503,8 @@ sycl::event inclusive_scan_base_step_striped(
                         slm_iscan_tmp[local_offset0 + m_wi] = local_iscan[m_wi];
                     }
 
-                    it.barrier(sycl::access::fence_space::local_space);
+                    sycl::group_barrier(it.get_group(),
+                                        sycl::memory_scope::work_group);
                 }
             }
 
diff --git a/dpnp/tensor/libtensor/include/kernels/linalg_functions/gemm.hpp b/dpnp/tensor/libtensor/include/kernels/linalg_functions/gemm.hpp
index 51e95b799bd0..e93a386ad785 100644
--- a/dpnp/tensor/libtensor/include/kernels/linalg_functions/gemm.hpp
+++ b/dpnp/tensor/libtensor/include/kernels/linalg_functions/gemm.hpp
@@ -653,7 +653,7 @@ class GemmBatchFunctorThreadK
             }
         }
 
-        it.barrier(sycl::access::fence_space::local_space);
+        sycl::group_barrier(it.get_group(), sycl::memory_scope::work_group);
 
         std::size_t t_shift = block_s * delta_k * n_wi;
         std::size_t global_s_offset = i * k + t_shift;
@@ -672,7 +672,7 @@ class GemmBatchFunctorThreadK
         std::size_t workspace_i_shift = local_i * delta_k;
         workspace[workspace_i_shift + local_s] = private_sum;
 
-        it.barrier(sycl::access::fence_space::local_space);
+        sycl::group_barrier(it.get_group(), sycl::memory_scope::work_group);
 
         if (local_s == 0 && i < n) {
             accV_t local_sum(workspace[workspace_i_shift]);
@@ -1063,7 +1063,7 @@ class GemmBatchFunctorThreadNM_vecm
                 }
             }
 
-            it.barrier(sycl::access::fence_space::local_space);
+            sycl::group_barrier(it.get_group(), sycl::memory_scope::work_group);
 
             const std::uint32_t lo_lhs_st_k = (wg_delta_n * wi_delta_n);
             const std::uint32_t lo_rhs_rk_k = (wg_delta_m * wi_delta_m_vecs);
@@ -1095,7 +1095,7 @@ class GemmBatchFunctorThreadNM_vecm
                 }
             }
 
-            it.barrier(sycl::access::fence_space::local_space);
+            sycl::group_barrier(it.get_group(), sycl::memory_scope::work_group);
         }
 
         if constexpr (m_vec_size == 1) {
@@ -1942,7 +1942,7 @@ class GemmBatchNoAtomicFunctorThreadNM
             }
         }
 
-        it.barrier(sycl::access::fence_space::local_space);
+        sycl::group_barrier(it.get_group(), sycl::memory_scope::work_group);
 
         i += local_i * wi_delta_n;
         j += local_j * wi_delta_m;
@@ -2115,7 +2115,7 @@ class GemmBatchNoAtomicFunctorThreadK
             }
         }
 
-        it.barrier(sycl::access::fence_space::local_space);
+        sycl::group_barrier(it.get_group(), sycl::memory_scope::work_group);
 
         std::size_t t_shift = block_s * delta_k * n_wi;
         std::size_t global_s_offset = i * k + t_shift;
@@ -2134,7 +2134,7 @@ class GemmBatchNoAtomicFunctorThreadK
         std::size_t workspace_i_shift = local_i * delta_k;
         workspace[workspace_i_shift + local_s] = private_sum;
 
-        it.barrier(sycl::access::fence_space::local_space);
+        sycl::group_barrier(it.get_group(), sycl::memory_scope::work_group);
 
         if (local_s == 0 && i < n) {
             accV_t local_sum(workspace[workspace_i_shift]);