diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 353cde70e8bf9..b59925c46c508 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -1326,8 +1326,6 @@ class __SYCL_EXPORT handler {
                   "Kernel argument cannot have a sycl::nd_item type in "
                   "sycl::parallel_for with sycl::range");
 
-#if defined(SYCL2020_CONFORMANT_APIS) ||                                       \
-    defined(__INTEL_PREVIEW_BREAKING_CHANGES)
     static_assert(std::is_convertible_v<item<Dims>, LambdaArgType> ||
                       std::is_convertible_v<item<Dims, false>, LambdaArgType>,
                   "sycl::parallel_for(sycl::range) kernel must have the "
@@ -1340,7 +1338,6 @@ class __SYCL_EXPORT handler {
          std::is_invocable_v<KernelType, RefLambdaArgType, kernel_handler>),
         "SYCL kernel lambda/functor has an unexpected signature, it should be "
         "invocable with sycl::item and optionally sycl::kernel_handler");
-#endif
 
     // TODO: Properties may change the kernel function, so in order to avoid
     //       conflicts they should be included in the name.
@@ -1432,19 +1429,11 @@ class __SYCL_EXPORT handler {
     verifyUsedKernelBundle(detail::KernelInfo<NameT>::getName());
     using LambdaArgType =
         sycl::detail::lambda_arg_type<KernelType, nd_item<Dims>>;
-#if defined(SYCL2020_CONFORMANT_APIS) ||                                       \
-    defined(__INTEL_PREVIEW_BREAKING_CHANGES)
     static_assert(
         std::is_convertible_v<sycl::nd_item<Dims>, LambdaArgType>,
         "Kernel argument of a sycl::parallel_for with sycl::nd_range "
         "must be either sycl::nd_item or be convertible from sycl::nd_item");
     using TransformedArgType = sycl::nd_item<Dims>;
-#else
-    // If user type is convertible from sycl::item/sycl::nd_item, use
-    // sycl::item/sycl::nd_item to transport item information
-    using TransformedArgType =
-        typename TransformUserItemType<Dims, LambdaArgType>::type;
-#endif
 
     (void)ExecutionRange;
     (void)Props;
diff --git a/sycl/test-e2e/Assert/assert_in_multiple_tus.hpp b/sycl/test-e2e/Assert/assert_in_multiple_tus.hpp
index d50af49f1ef4e..585d75dfde7a0 100644
--- a/sycl/test-e2e/Assert/assert_in_multiple_tus.hpp
+++ b/sycl/test-e2e/Assert/assert_in_multiple_tus.hpp
@@ -32,11 +32,12 @@ void enqueueKernel_1_fromFile1(queue *Q) {
 
     CGH.parallel_for<class Kernel_1>(
         sycl::nd_range(Buf.get_range(), sycl::range<1>(4)),
-        [=](sycl::id<1> wiID) {
+        [=](sycl::nd_item<1> ndi) {
+          auto gid = ndi.get_global_id(0);
           int X = 0;
-          if (wiID == 5)
+          if (gid == 5)
             X = checkFunction();
-          Acc[wiID] = X;
+          Acc[gid] = X;
         });
   });
 }
diff --git a/sycl/test-e2e/DiscardEvents/discard_events_accessors.cpp b/sycl/test-e2e/DiscardEvents/discard_events_accessors.cpp
index fa9af30e33a05..68c68d1c63940 100644
--- a/sycl/test-e2e/DiscardEvents/discard_events_accessors.cpp
+++ b/sycl/test-e2e/DiscardEvents/discard_events_accessors.cpp
@@ -59,8 +59,8 @@ int main(int Argc, const char *Argv[]) {
       sycl::local_accessor<int, 1> LocalAcc(LocalMemSize, CGH);
 
       CGH.parallel_for<class kernel_using_local_memory>(
-          NDRange, [=](sycl::item<1> itemID) {
-            size_t i = itemID.get_id(0);
+          NDRange, [=](sycl::nd_item<1> ndi) {
+            size_t i = ndi.get_global_id(0);
             int *Ptr = LocalAcc.get_pointer();
             Ptr[i] = i + 5;
             Harray[i] = Ptr[i] + 5;
diff --git a/sycl/test-e2e/ESIMD/bfn.cpp b/sycl/test-e2e/ESIMD/bfn.cpp
index 6849399453927..208ae7ccbbfb7 100644
--- a/sycl/test-e2e/ESIMD/bfn.cpp
+++ b/sycl/test-e2e/ESIMD/bfn.cpp
@@ -107,8 +107,9 @@ struct DeviceFunc {
   DeviceFunc(const T *In0, const T *In1, const T *In2, T *Out)
       : In0(In0), In1(In1), In2(In2), Out(Out) {}
 
-  void operator()(id<1> I) const SYCL_ESIMD_KERNEL {
-    unsigned int Offset = I * N;
+  void operator()(nd_item<1> ndi) const SYCL_ESIMD_KERNEL {
+    auto gid = ndi.get_global_id(0);
+    unsigned int Offset = gid * N;
     esimd::simd<T, N> V0;
     esimd::simd<T, N> V1;
     esimd::simd<T, N> V2;
@@ -116,7 +117,7 @@ struct DeviceFunc {
     V1.copy_from(In1 + Offset);
     V2.copy_from(In2 + Offset);
 
-    if (I.get(0) % 2 == 0) {
+    if (gid % 2 == 0) {
       for (int J = 0; J < N; J++) {
         Kernel<T, N, Op, AllSca> DevF{};
         T Val0 = V0[J];
diff --git a/sycl/test-e2e/ESIMD/ext_math.cpp b/sycl/test-e2e/ESIMD/ext_math.cpp
index f71ffeafb9762..2ee43592ed8ba 100644
--- a/sycl/test-e2e/ESIMD/ext_math.cpp
+++ b/sycl/test-e2e/ESIMD/ext_math.cpp
@@ -238,12 +238,13 @@ struct UnaryDeviceFunc {
 
   UnaryDeviceFunc(AccIn &In, AccOut &Out) : In(In), Out(Out) {}
 
-  void operator()(id<1> I) const SYCL_ESIMD_KERNEL {
-    unsigned int Offset = I * N * sizeof(T);
+  void operator()(nd_item<1> ndi) const SYCL_ESIMD_KERNEL {
+    auto gid = ndi.get_global_id(0);
+    unsigned int Offset = gid * N * sizeof(T);
     esimd::simd<T, N> Vx;
     Vx.copy_from(In, Offset);
 
-    if (I.get(0) % 2 == 0) {
+    if (gid % 2 == 0) {
       for (int J = 0; J < N; J++) {
         Kernel<T, N, Op, AllSca> DevF{};
         T Val = Vx[J];
@@ -269,13 +270,14 @@ struct BinaryDeviceFunc {
   BinaryDeviceFunc(AccIn &In1, AccIn &In2, AccOut &Out)
       : In1(In1), In2(In2), Out(Out) {}
 
-  void operator()(id<1> I) const SYCL_ESIMD_KERNEL {
-    unsigned int Offset = I * N * sizeof(T);
+  void operator()(nd_item<1> ndi) const SYCL_ESIMD_KERNEL {
+    auto gid = ndi.get_global_id(0);
+    unsigned int Offset = gid * N * sizeof(T);
     esimd::simd<T, N> V1(In1, Offset);
     esimd::simd<T, N> V2(In2, Offset);
     esimd::simd<T, N> V;
 
-    if (I.get(0) % 2 == 0) {
+    if (gid % 2 == 0) {
       int Ind = 0;
       {
         Kernel<T, N, Op, AllSca> DevF{};
diff --git a/sycl/test-e2e/ESIMD/local_accessor_gather_scatter.cpp b/sycl/test-e2e/ESIMD/local_accessor_gather_scatter.cpp
index a3262ae764dc9..9b9ecd2fef177 100644
--- a/sycl/test-e2e/ESIMD/local_accessor_gather_scatter.cpp
+++ b/sycl/test-e2e/ESIMD/local_accessor_gather_scatter.cpp
@@ -50,7 +50,7 @@ template <typename T, unsigned VL, unsigned STRIDE> bool test(queue q) {
     q.submit([&](handler &cgh) {
        auto acc = buf.template get_access<access::mode::read_write>(cgh);
        auto LocalAcc = local_accessor<T, 1>(size * STRIDE, cgh);
-       cgh.parallel_for(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL {
+       cgh.parallel_for(glob_range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
          using namespace sycl::ext::intel::esimd;
          simd<T, VL> valsIn;
          valsIn.copy_from(acc, 0);
diff --git a/sycl/test-e2e/ESIMD/local_accessor_gather_scatter_rgba.cpp b/sycl/test-e2e/ESIMD/local_accessor_gather_scatter_rgba.cpp
index 7917899a947fa..5153ef0eae76c 100644
--- a/sycl/test-e2e/ESIMD/local_accessor_gather_scatter_rgba.cpp
+++ b/sycl/test-e2e/ESIMD/local_accessor_gather_scatter_rgba.cpp
@@ -89,7 +89,7 @@ template <typename T, unsigned VL, auto CH_MASK> bool test(queue q) {
        auto OutAcc = OutBuf.template get_access<access::mode::read_write>(cgh);
        auto LocalAcc = local_accessor<T, 1>(VL * NUM_RGBA_CHANNELS, cgh);
 
-       cgh.parallel_for(Range, [=](id<1> i) SYCL_ESIMD_KERNEL {
+       cgh.parallel_for(Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
          using namespace sycl::ext::intel::esimd;
          constexpr int numChannels = get_num_channels_enabled(CH_MASK);
 
@@ -135,7 +135,7 @@ template <typename T, unsigned VL, auto CH_MASK> bool test(queue q) {
                -1;
          }
 
-         uint32_t global_offset = i * VL * NUM_RGBA_CHANNELS;
+         uint32_t global_offset = ndi.get_global_id(0) * VL * NUM_RGBA_CHANNELS;
          valsOut.copy_to(OutAcc, global_offset);
        });
      }).wait();
diff --git a/sycl/test-e2e/ESIMD/lsc/atomic_smoke.cpp b/sycl/test-e2e/ESIMD/lsc/atomic_smoke.cpp
index bb653f7246dd1..c234e86e31a5a 100644
--- a/sycl/test-e2e/ESIMD/lsc/atomic_smoke.cpp
+++ b/sycl/test-e2e/ESIMD/lsc/atomic_smoke.cpp
@@ -212,8 +212,8 @@ bool test(queue q, const Config &cfg) {
   try {
     auto e = q.submit([&](handler &cgh) {
       cgh.parallel_for<TestID<T, N, ImplF>>(
-          rng, [=](id<1> ii) SYCL_ESIMD_KERNEL {
-            int i = ii;
+          rng, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
+            int i = ndi.get_global_id(0);
 #ifndef USE_SCALAR_OFFSET
             simd<Toffset, N> offsets(cfg.start_ind * sizeof(T),
                                      cfg.stride * sizeof(T));
@@ -332,8 +332,8 @@ bool test(queue q, const Config &cfg) {
     auto e = q.submit([&](handler &cgh) {
       auto accessor = buf.template get_access<access::mode::read_write>(cgh);
       cgh.parallel_for<TestID<T, N, ImplF>>(
-          rng, [=](id<1> ii) SYCL_ESIMD_KERNEL {
-            int i = ii;
+          rng, [=](nd_item<1> gid) SYCL_ESIMD_KERNEL {
+            int i = gid.get_global_id(0);
 #ifndef USE_SCALAR_OFFSET
             simd<Toffset, N> offsets(start * sizeof(T), stride * sizeof(T));
 #else
diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp
index 02c8fb1bfe9a2..a75bed514b497 100644
--- a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp
+++ b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp
@@ -52,7 +52,7 @@ template <typename T, unsigned VL, unsigned STRIDE> bool test(queue q) {
     q.submit([&](handler &cgh) {
        auto acc = buf.template get_access<access::mode::read_write>(cgh);
        auto LocalAcc = local_accessor<T, 1>(size * STRIDE, cgh);
-       cgh.parallel_for(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL {
+       cgh.parallel_for(glob_range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
          using namespace sycl::ext::intel::esimd;
          simd<T, VL> valsIn;
          valsIn.copy_from(acc, 0);
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp b/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp
index 62315c4d58727..985f44c0e6954 100644
--- a/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp
@@ -153,8 +153,8 @@ bool test_usm(queue q, const Config &cfg) {
 
   try {
     auto e = q.submit([&](handler &cgh) {
-      cgh.parallel_for(rng, [=](id<1> ii) SYCL_ESIMD_KERNEL {
-        int i = ii;
+      cgh.parallel_for(rng, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
+        int i = ndi.get_global_id(0);
         simd<Toffset, N> offsets(cfg.start_ind * sizeof(T),
                                  cfg.stride * sizeof(T));
         simd_mask<N> m = 1;
@@ -287,8 +287,8 @@ bool test_acc(queue q, const Config &cfg) {
       auto e = q.submit([&](handler &cgh) {
         auto arr_acc =
             arr_buf.template get_access<access::mode::read_write>(cgh);
-        cgh.parallel_for(rng, [=](id<1> ii) SYCL_ESIMD_KERNEL {
-          int i = ii;
+        cgh.parallel_for(rng, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
+          int i = ndi.get_global_id(0);
           simd<Toffset, N> offsets(cfg.start_ind * sizeof(T),
                                    cfg.stride * sizeof(T));
           simd_mask<N> m = 1;
diff --git a/sycl/test-e2e/KernelFusion/abort_fusion.cpp b/sycl/test-e2e/KernelFusion/abort_fusion.cpp
index e50d701a9ee7d..ba8294e5b49d8 100644
--- a/sycl/test-e2e/KernelFusion/abort_fusion.cpp
+++ b/sycl/test-e2e/KernelFusion/abort_fusion.cpp
@@ -25,12 +25,16 @@ template <int Dimensions> size_t getSize(nd_range<Dimensions> r) {
   return r.get_global_range().size();
 }
 
+template <int N> auto global_linear_id(sycl::nd_item<N> ndi) {
+  return ndi.get_global_linear_id();
+}
+template <int N> auto global_linear_id(sycl::item<N> i) {
+  return i.get_linear_id();
+}
+
 template <typename Kernel1Name, typename Kernel2Name, typename Range1,
           typename Range2>
 void performFusion(queue &q, Range1 R1, Range2 R2) {
-  using IndexTy1 = item<Range1::dimensions>;
-  using IndexTy2 = item<Range2::dimensions>;
-
   int in[dataSize], tmp[dataSize], out[dataSize];
 
   for (size_t i = 0; i < dataSize; ++i) {
@@ -51,8 +55,8 @@ void performFusion(queue &q, Range1 R1, Range2 R2) {
     q.submit([&](handler &cgh) {
       auto accIn = bIn.get_access(cgh);
       auto accTmp = bTmp.get_access(cgh);
-      cgh.parallel_for<Kernel1Name>(R1, [=](IndexTy1 i) {
-        size_t j = i.get_linear_id();
+      cgh.parallel_for<Kernel1Name>(R1, [=](auto i) {
+        size_t j = global_linear_id(i);
         accTmp[j] = accIn[j] + 5;
       });
     });
@@ -60,8 +64,8 @@ void performFusion(queue &q, Range1 R1, Range2 R2) {
     q.submit([&](handler &cgh) {
       auto accTmp = bTmp.get_access(cgh);
       auto accOut = bOut.get_access(cgh);
-      cgh.parallel_for<Kernel2Name>(R2, [=](IndexTy2 i) {
-        size_t j = i.get_linear_id();
+      cgh.parallel_for<Kernel2Name>(R2, [=](auto i) {
+        size_t j = global_linear_id(i);
         accOut[j] = accTmp[j] * 2;
       });
     });
diff --git a/sycl/test-e2e/KernelFusion/abort_internalization.cpp b/sycl/test-e2e/KernelFusion/abort_internalization.cpp
index 3dde5189529c8..f1838e96fde0f 100644
--- a/sycl/test-e2e/KernelFusion/abort_internalization.cpp
+++ b/sycl/test-e2e/KernelFusion/abort_internalization.cpp
@@ -49,8 +49,10 @@ void performFusion(queue &q, Internalization intKernel1,
 
       if (localSizeKernel1 > 0) {
         cgh.parallel_for<class Kernel1>(
-            nd_range<1>{{dataSize}, {localSizeKernel1}},
-            [=](id<1> i) { accTmp[i] = accIn[i] + 5; });
+            nd_range<1>{{dataSize}, {localSizeKernel1}}, [=](nd_item<1> ndi) {
+              auto i = ndi.get_global_id(0);
+              accTmp[i] = accIn[i] + 5;
+            });
       } else {
         cgh.parallel_for<class KernelOne>(
             dataSize, [=](id<1> i) { accTmp[i] = accIn[i] + 5; });
@@ -70,8 +72,10 @@ void performFusion(queue &q, Internalization intKernel1,
       auto accOut = bOut.get_access(cgh);
       if (localSizeKernel2 > 0) {
         cgh.parallel_for<class Kernel2>(
-            nd_range<1>{{dataSize}, {localSizeKernel2}},
-            [=](id<1> i) { accOut[i] = accTmp[i] * 2; });
+            nd_range<1>{{dataSize}, {localSizeKernel2}}, [=](nd_item<1> ndi) {
+              auto i = ndi.get_global_id(0);
+              accOut[i] = accTmp[i] * 2;
+            });
       } else {
         cgh.parallel_for<class KernelTwo>(
             dataSize, [=](id<1> i) { accOut[i] = accTmp[i] * 2; });
diff --git a/sycl/test-e2e/KernelFusion/diamond_shape_local.cpp b/sycl/test-e2e/KernelFusion/diamond_shape_local.cpp
index bd16ab8e3a295..ce5df8fbc20cd 100644
--- a/sycl/test-e2e/KernelFusion/diamond_shape_local.cpp
+++ b/sycl/test-e2e/KernelFusion/diamond_shape_local.cpp
@@ -16,7 +16,10 @@ struct AddKernel {
   accessor<int, 1> accIn2;
   accessor<int, 1> accOut;
 
-  void operator()(id<1> i) const { accOut[i] = accIn1[i] + accIn2[i]; }
+  void operator()(nd_item<1> ndi) const {
+    auto i = ndi.get_global_id(0);
+    accOut[i] = accIn1[i] + accIn2[i];
+  }
 };
 
 int main() {
@@ -71,17 +74,21 @@ int main() {
       auto accTmp1 = bTmp1.get_access(cgh);
       auto accIn3 = bIn3.get_access(cgh);
       auto accTmp2 = bTmp2.get_access(cgh);
-      cgh.parallel_for<class KernelOne>(
-          nd_range<1>{{dataSize}, {16}},
-          [=](id<1> i) { accTmp2[i] = accTmp1[i] * accIn3[i]; });
+      cgh.parallel_for<class KernelOne>(nd_range<1>{{dataSize}, {16}},
+                                        [=](nd_item<1> ndi) {
+                                          auto i = ndi.get_global_id(0);
+                                          accTmp2[i] = accTmp1[i] * accIn3[i];
+                                        });
     });
 
     q.submit([&](handler &cgh) {
       auto accTmp1 = bTmp1.get_access(cgh);
       auto accTmp3 = bTmp3.get_access(cgh);
-      cgh.parallel_for<class KernelTwo>(
-          nd_range<1>{{dataSize}, {16}},
-          [=](id<1> i) { accTmp3[i] = accTmp1[i] * 5; });
+      cgh.parallel_for<class KernelTwo>(nd_range<1>{{dataSize}, {16}},
+                                        [=](nd_item<1> ndi) {
+                                          auto i = ndi.get_global_id(0);
+                                          accTmp3[i] = accTmp1[i] * 5;
+                                        });
     });
 
     q.submit([&](handler &cgh) {
diff --git a/sycl/test-e2e/KernelFusion/existing_local_accessor.cpp b/sycl/test-e2e/KernelFusion/existing_local_accessor.cpp
index b41c7caf86b9f..5b168cfca9332 100644
--- a/sycl/test-e2e/KernelFusion/existing_local_accessor.cpp
+++ b/sycl/test-e2e/KernelFusion/existing_local_accessor.cpp
@@ -56,9 +56,11 @@ int main() {
           cgh, sycl::ext::codeplay::experimental::property::promote_local{});
       auto accIn3 = bIn3.get_access(cgh);
       auto accOut = bOut.get_access(cgh);
-      cgh.parallel_for<class KernelTwo>(
-          nd_range<1>{{dataSize}, {16}},
-          [=](id<1> i) { accOut[i] = accTmp[i] * accIn3[i]; });
+      cgh.parallel_for<class KernelTwo>(nd_range<1>{{dataSize}, {16}},
+                                        [=](nd_item<1> ndi) {
+                                          auto i = ndi.get_global_id(0);
+                                          accOut[i] = accTmp[i] * accIn3[i];
+                                        });
     });
 
     fw.complete_fusion({ext::codeplay::experimental::property::no_barriers{}});
diff --git a/sycl/test-e2e/KernelFusion/internalize_array_wrapper_local.cpp b/sycl/test-e2e/KernelFusion/internalize_array_wrapper_local.cpp
index bb3607b79abc3..e019c720a3156 100644
--- a/sycl/test-e2e/KernelFusion/internalize_array_wrapper_local.cpp
+++ b/sycl/test-e2e/KernelFusion/internalize_array_wrapper_local.cpp
@@ -84,7 +84,8 @@ int main() {
       auto accTmp = bTmp.get_access(
           cgh, sycl::ext::codeplay::experimental::property::promote_local{});
       cgh.parallel_for<class KernelOne>(
-          nd_range<1>{{dataSize}, {4}}, [=](id<1> id) {
+          nd_range<1>{{dataSize}, {4}}, [=](nd_item<1> ndi) {
+            auto id = ndi.get_global_id();
             const auto &accIn1Wrapp = accIn1[id];
             const auto &accIn2Wrapp = accIn2[id];
             auto &accTmpWrapp = accTmp[id];
@@ -105,7 +106,8 @@ int main() {
       auto accIn3 = bIn3.get_access(cgh);
       auto accOut = bOut.get_access(cgh);
       cgh.parallel_for<class KernelTwo>(
-          nd_range<1>{{dataSize}, {4}}, [=](id<1> id) {
+          nd_range<1>{{dataSize}, {4}}, [=](nd_item<1> ndi) {
+            auto id = ndi.get_global_id();
             const auto &tmpWrapp = accTmp[id];
             const auto &accIn3Wrapp = accIn3[id];
             auto &accOutWrapp = accOut[id];
diff --git a/sycl/test-e2e/KernelFusion/local_internalization.cpp b/sycl/test-e2e/KernelFusion/local_internalization.cpp
index 1d84291f9d3f5..c18ab5cf92d36 100644
--- a/sycl/test-e2e/KernelFusion/local_internalization.cpp
+++ b/sycl/test-e2e/KernelFusion/local_internalization.cpp
@@ -41,9 +41,11 @@ int main() {
       auto accIn2 = bIn2.get_access(cgh);
       auto accTmp = bTmp.get_access(
           cgh, sycl::ext::codeplay::experimental::property::promote_local{});
-      cgh.parallel_for<class KernelOne>(
-          nd_range<1>{{dataSize}, {16}},
-          [=](id<1> i) { accTmp[i] = accIn1[i] + accIn2[i]; });
+      cgh.parallel_for<class KernelOne>(nd_range<1>{{dataSize}, {16}},
+                                        [=](nd_item<1> ndi) {
+                                          auto i = ndi.get_global_id(0);
+                                          accTmp[i] = accIn1[i] + accIn2[i];
+                                        });
     });
 
     q.submit([&](handler &cgh) {
@@ -51,9 +53,11 @@ int main() {
           cgh, sycl::ext::codeplay::experimental::property::promote_local{});
       auto accIn3 = bIn3.get_access(cgh);
       auto accOut = bOut.get_access(cgh);
-      cgh.parallel_for<class KernelTwo>(
-          nd_range<1>{{dataSize}, {16}},
-          [=](id<1> i) { accOut[i] = accTmp[i] * accIn3[i]; });
+      cgh.parallel_for<class KernelTwo>(nd_range<1>{{dataSize}, {16}},
+                                        [=](nd_item<1> ndi) {
+                                          auto i = ndi.get_global_id(0);
+                                          accOut[i] = accTmp[i] * accIn3[i];
+                                        });
     });
 
     fw.complete_fusion({ext::codeplay::experimental::property::no_barriers{}});
diff --git a/sycl/test-e2e/KernelFusion/non_unit_local_size.cpp b/sycl/test-e2e/KernelFusion/non_unit_local_size.cpp
index 82345f475cb32..120c02bb57418 100644
--- a/sycl/test-e2e/KernelFusion/non_unit_local_size.cpp
+++ b/sycl/test-e2e/KernelFusion/non_unit_local_size.cpp
@@ -42,8 +42,8 @@ int main() {
       auto accTmp = bTmp.get_access(
           cgh, sycl::ext::codeplay::experimental::property::promote_local{});
       cgh.parallel_for<class KernelOne>(
-          nd_range<1>{{128}, {8}}, [=](item<1> i) {
-            auto baseOffset = i.get_linear_id() * 4;
+          nd_range<1>{{128}, {8}}, [=](nd_item<1> ndi) {
+            auto baseOffset = ndi.get_global_linear_id() * 4;
             for (size_t j = 0; j < 4; ++j) {
               accTmp[baseOffset + j] =
                   accIn1[baseOffset + j] + accIn2[baseOffset + j];
@@ -57,8 +57,8 @@ int main() {
       auto accIn3 = bIn3.get_access(cgh);
       auto accOut = bOut.get_access(cgh);
       cgh.parallel_for<class KernelTwo>(
-          nd_range<1>{{128}, {8}}, [=](item<1> i) {
-            auto baseOffset = i.get_linear_id() * 4;
+          nd_range<1>{{128}, {8}}, [=](nd_item<1> ndi) {
+            auto baseOffset = ndi.get_global_linear_id() * 4;
             for (size_t j = 0; j < 4; ++j) {
               accOut[baseOffset + j] =
                   accTmp[baseOffset + j] * accIn3[baseOffset + j];
diff --git a/sycl/test-e2e/KernelFusion/struct_with_array.cpp b/sycl/test-e2e/KernelFusion/struct_with_array.cpp
index 7ccec0ee7fce8..5f35d0d2e7c23 100644
--- a/sycl/test-e2e/KernelFusion/struct_with_array.cpp
+++ b/sycl/test-e2e/KernelFusion/struct_with_array.cpp
@@ -53,9 +53,11 @@ int main() {
       auto accIn2 = bIn2.get_access(cgh);
       auto accTmp = bTmp.get_access(
           cgh, sycl::ext::codeplay::experimental::property::promote_private{});
-      cgh.parallel_for<class KernelOne>(
-          nd_range<1>{{dataSize}, {8}},
-          [=](id<1> i) { accTmp[i] = accIn1[i] + accIn2[i]; });
+      cgh.parallel_for<class KernelOne>(nd_range<1>{{dataSize}, {8}},
+                                        [=](nd_item<1> ndi) {
+                                          auto i = ndi.get_global_id(0);
+                                          accTmp[i] = accIn1[i] + accIn2[i];
+                                        });
     });
 
     q.submit([&](handler &cgh) {
diff --git a/sycl/test-e2e/Tracing/code_location_queue_parallel_for.cpp b/sycl/test-e2e/Tracing/code_location_queue_parallel_for.cpp
index b0d4249f91723..360d47d49a5b8 100644
--- a/sycl/test-e2e/Tracing/code_location_queue_parallel_for.cpp
+++ b/sycl/test-e2e/Tracing/code_location_queue_parallel_for.cpp
@@ -18,7 +18,7 @@ int main() {
     // CHECK: code_location_queue_parallel_for.cpp:[[# @LINE + 3 ]] E2ETestKernel
     Queue.parallel_for<class E2ETestKernel>(
         sycl::nd_range<1>{MaxWISizes.get(0), 2 * MaxWISizes.get(0)},
-        [](sycl::id<1> idx) {});
+        [](sycl::nd_item<1>) {});
   } catch (...) {
     ExceptionCaught = true;
   }
diff --git a/sycl/test-e2e/XPTI/buffer/accessors.cpp b/sycl/test-e2e/XPTI/buffer/accessors.cpp
index 087df8854fe2e..094aa08cc7285 100644
--- a/sycl/test-e2e/XPTI/buffer/accessors.cpp
+++ b/sycl/test-e2e/XPTI/buffer/accessors.cpp
@@ -36,7 +36,7 @@ int main() {
     auto A5 = Buf.get_access<mode::discard_read_write, target::device>(cgh);
     // CHECK: {{[0-9]+}}|Construct accessor|[[BUFFERID]]|[[ACCID6:.*]]|2014|1029|{{.*}}accessors.cpp:[[# @LINE + 1]]:15
     auto A6 = Buf.get_access<mode::atomic>(cgh);
-    cgh.parallel_for<class FillBuffer>(NDRange, [=](sycl::id<1> WIid) {
+    cgh.parallel_for<class FillBuffer>(NDRange, [=](sycl::nd_item<1>) {
       (void)A1;
       (void)A2;
       (void)A3;
diff --git a/sycl/test-e2e/XPTI/kernel/basic.cpp b/sycl/test-e2e/XPTI/kernel/basic.cpp
index 45ce046b5a07a..fe433ecd438fb 100644
--- a/sycl/test-e2e/XPTI/kernel/basic.cpp
+++ b/sycl/test-e2e/XPTI/kernel/basic.cpp
@@ -59,18 +59,19 @@ int main() {
         auto A1 = Buf.get_access<mode::read_write>(cgh);
         // CHECK: {{[0-9]+}}|Construct accessor|0x0|[[ACCID2:.*]]|2016|1026|{{.*}}.cpp:[[# @LINE + 1]]:38
         sycl::local_accessor<int, 1> A2(Range, cgh);
-        cgh.parallel_for<class FillBuffer>(NDRange, [=](sycl::id<1> WIid) {
+        cgh.parallel_for<class FillBuffer>(NDRange, [=](sycl::nd_item<1> ndi) {
+          auto gid = ndi.get_global_id(0);
           // CHECK-OPT: arg0 : {1, {{[0-9,a-f,x]+}}, 2, 0}
           int h = Val;
           // CHECK-OPT: arg1 : {1, {{.*}}0, 20, 1}
-          A2[WIid[0]] = h;
+          A2[gid] = h;
           // CHECK-OPT: arg2 : {0, [[ACCID1]], 4062, 2}
           // CHECK-OPT: arg3 : {1, [[ACCID1]], 8, 3}
-          A1[WIid[0]] = A2[WIid[0]];
+          A1[gid] = A2[gid];
           // CHECK-OPT: arg4 : {3, {{.*}}, 8, 4}
-          PtrDevice[WIid[0]] = WIid[0];
+          PtrDevice[gid] = gid;
           // CHECK-OPT: arg5 : {3, {{.*}}, 8, 5}
-          PtrShared[WIid[0]] = PtrDevice[WIid[0]];
+          PtrShared[gid] = PtrDevice[gid];
         });
       })
       .wait();
diff --git a/sycl/test/check_device_code/native_cpu/native_cpu_builtins.cpp b/sycl/test/check_device_code/native_cpu/native_cpu_builtins.cpp
index 783fd6a30a7e1..d4b088ac3c0f7 100644
--- a/sycl/test/check_device_code/native_cpu/native_cpu_builtins.cpp
+++ b/sycl/test/check_device_code/native_cpu/native_cpu_builtins.cpp
@@ -44,15 +44,17 @@ int main() {
                                    1,
                                });
   deviceQueue.submit([&](sycl::handler &h) {
-    h.parallel_for<Test2>(r2, [=](sycl::id<2> id) { acc[id[1]] = 42; });
+    h.parallel_for<Test2>(
+        r2, [=](sycl::nd_item<2> ndi) { acc[ndi.get_global_id(1)] = 42; });
     // CHECK: @_ZTS5Test2.NativeCPUKernel(ptr {{.*}}%0, ptr {{.*}}%1, ptr addrspace(1) %2)
     // CHECK: call{{.*}}__dpcpp_nativecpu_get_global_id(i32 1, ptr addrspace(1) %2)
     // CHECK: call{{.*}}__dpcpp_nativecpu_get_global_id(i32 0, ptr addrspace(1) %2)
   });
   sycl::nd_range<3> r3({1, 1, 1}, {1, 1, 1});
   deviceQueue.submit([&](sycl::handler &h) {
-    h.parallel_for<Test3>(
-        r3, [=](sycl::item<3> item) { acc[item[2]] = item.get_range(0); });
+    h.parallel_for<Test3>(r3, [=](sycl::nd_item<3> ndi) {
+      acc[ndi.get_global_id(2)] = ndi.get_global_range(0);
+    });
     // CHECK: @_ZTS5Test3.NativeCPUKernel(ptr {{.*}}%0, ptr {{.*}}%1, ptr addrspace(1) %2)
     // CHECK-DAG: call{{.*}}__dpcpp_nativecpu_get_global_range(i32 2, ptr addrspace(1) %2)
     // CHECK-DAG: call{{.*}}__dpcpp_nativecpu_get_global_range(i32 1, ptr addrspace(1) %2)
diff --git a/sycl/unittests/Extensions/CommandGraph/Update.cpp b/sycl/unittests/Extensions/CommandGraph/Update.cpp
index 92246fb83678d..0f17a12b4680f 100644
--- a/sycl/unittests/Extensions/CommandGraph/Update.cpp
+++ b/sycl/unittests/Extensions/CommandGraph/Update.cpp
@@ -128,7 +128,7 @@ TEST_F(CommandGraphTest, UpdateRangeErrors) {
   nd_range<1> NDRange{range{128}, range{32}};
   range<1> Range{128};
   auto NodeNDRange = Graph.add([&](sycl::handler &cgh) {
-    cgh.parallel_for<TestKernel<>>(NDRange, [](item<1>) {});
+    cgh.parallel_for<TestKernel<>>(NDRange, [](nd_item<1>) {});
   });
 
   // OK