diff --git a/sycl/include/CL/sycl.hpp b/sycl/include/CL/sycl.hpp
index 659dee0ed6aa3..383725dcc4b88 100644
--- a/sycl/include/CL/sycl.hpp
+++ b/sycl/include/CL/sycl.hpp
@@ -35,6 +35,7 @@
 #include <CL/sycl/item.hpp>
 #include <CL/sycl/kernel.hpp>
 #include <CL/sycl/kernel_bundle.hpp>
+#include <CL/sycl/kernel_handler.hpp>
 #include <CL/sycl/marray.hpp>
 #include <CL/sycl/multi_ptr.hpp>
 #include <CL/sycl/nd_item.hpp>
@@ -48,6 +49,7 @@
 #include <CL/sycl/range.hpp>
 #include <CL/sycl/reduction.hpp>
 #include <CL/sycl/sampler.hpp>
+#include <CL/sycl/specialization_id.hpp>
 #include <CL/sycl/stream.hpp>
 #include <CL/sycl/types.hpp>
 #include <CL/sycl/usm.hpp>
diff --git a/sycl/include/CL/sycl/detail/cg_types.hpp b/sycl/include/CL/sycl/detail/cg_types.hpp
index f8acc96550fda..5faa502102ba3 100644
--- a/sycl/include/CL/sycl/detail/cg_types.hpp
+++ b/sycl/include/CL/sycl/detail/cg_types.hpp
@@ -15,6 +15,7 @@
 #include <CL/sycl/interop_handle.hpp>
 #include <CL/sycl/interop_handler.hpp>
 #include <CL/sycl/kernel.hpp>
+#include <CL/sycl/kernel_handler.hpp>
 #include <CL/sycl/nd_item.hpp>
 #include <CL/sycl/range.hpp>
 
@@ -122,6 +123,97 @@ class NDRDescT {
   size_t Dims;
 };
 
+template <typename, typename T> struct check_fn_signature {
+  static_assert(std::integral_constant<T, false>::value,
+                "Second template parameter is required to be of function type");
+};
+
+template <typename F, typename RetT, typename... Args>
+struct check_fn_signature<F, RetT(Args...)> {
+private:
+  template <typename T>
+  static constexpr auto check(T *) -> typename std::is_same<
+      decltype(std::declval<T>().operator()(std::declval<Args>()...)),
+      RetT>::type;
+
+  template <typename> static constexpr std::false_type check(...);
+
+  using type = decltype(check<F>(0));
+
+public:
+  static constexpr bool value = type::value;
+};
+
+template <typename F, typename... Args>
+static constexpr bool check_kernel_lambda_takes_args() {
+  return check_fn_signature<std::remove_reference_t<F>, void(Args...)>::value;
+}
+
+// isKernelLambdaCallableWithKernelHandlerImpl checks if LambdaArgType is void
+// (e.g., in single_task), and based on that, calls
+// check_kernel_lambda_takes_args with proper set of arguments. Also this type
+// trait workarounds compilation error which happens only with msvc.
+
+template <typename KernelType, typename LambdaArgType,
+          typename std::enable_if_t<std::is_same<LambdaArgType, void>::value>
+              * = nullptr>
+constexpr bool isKernelLambdaCallableWithKernelHandlerImpl() {
+  return check_kernel_lambda_takes_args<KernelType, kernel_handler>();
+}
+
+template <typename KernelType, typename LambdaArgType,
+          typename std::enable_if_t<!std::is_same<LambdaArgType, void>::value>
+              * = nullptr>
+constexpr bool isKernelLambdaCallableWithKernelHandlerImpl() {
+  return check_kernel_lambda_takes_args<KernelType, LambdaArgType,
+                                        kernel_handler>();
+}
+
+// Type traits to find out if kernal lambda has kernel_handler argument
+
+template <typename KernelType>
+constexpr bool isKernelLambdaCallableWithKernelHandler() {
+  return check_kernel_lambda_takes_args<KernelType, kernel_handler>();
+}
+
+template <typename KernelType, typename LambdaArgType>
+constexpr bool isKernelLambdaCallableWithKernelHandler() {
+  return isKernelLambdaCallableWithKernelHandlerImpl<KernelType,
+                                                     LambdaArgType>();
+}
+
+// Helpers for running kernel lambda on the host device
+
+template <typename KernelType,
+          typename std::enable_if_t<isKernelLambdaCallableWithKernelHandler<
+              KernelType>()> * = nullptr>
+constexpr void runKernelWithoutArg(KernelType KernelName) {
+  kernel_handler KH;
+  KernelName(KH);
+}
+
+template <typename KernelType,
+          typename std::enable_if_t<!isKernelLambdaCallableWithKernelHandler<
+              KernelType>()> * = nullptr>
+constexpr void runKernelWithoutArg(KernelType KernelName) {
+  KernelName();
+}
+
+template <typename ArgType, typename KernelType,
+          typename std::enable_if_t<isKernelLambdaCallableWithKernelHandler<
+              KernelType, ArgType>()> * = nullptr>
+constexpr void runKernelWithArg(KernelType KernelName, ArgType Arg) {
+  kernel_handler KH;
+  KernelName(Arg, KH);
+}
+
+template <typename ArgType, typename KernelType,
+          typename std::enable_if_t<!isKernelLambdaCallableWithKernelHandler<
+              KernelType, ArgType>()> * = nullptr>
+constexpr void runKernelWithArg(KernelType KernelName, ArgType Arg) {
+  KernelName(Arg);
+}
+
 // The pure virtual class aimed to store lambda/functors of any type.
 class HostKernelBase {
 public:
@@ -197,7 +289,7 @@ class HostKernel : public HostKernelBase {
   template <class ArgT = KernelArgType>
   typename detail::enable_if_t<std::is_same<ArgT, void>::value>
   runOnHost(const NDRDescT &) {
-    MKernel();
+    runKernelWithoutArg(MKernel);
   }
 
   template <class ArgT = KernelArgType>
@@ -218,18 +310,18 @@ class HostKernel : public HostKernelBase {
       UpperBound[I] = Range[I] + Offset[I];
     }
 
-    detail::NDLoop<Dims>::iterate(/*LowerBound=*/Offset, Stride, UpperBound,
-                                  [&](const sycl::id<Dims> &ID) {
-                                    sycl::item<Dims, /*Offset=*/true> Item =
-                                        IDBuilder::createItem<Dims, true>(
-                                            Range, ID, Offset);
-
-                                    if (StoreLocation) {
-                                      store_id(&ID);
-                                      store_item(&Item);
-                                    }
-                                    MKernel(ID);
-                                  });
+    detail::NDLoop<Dims>::iterate(
+        /*LowerBound=*/Offset, Stride, UpperBound,
+        [&](const sycl::id<Dims> &ID) {
+          sycl::item<Dims, /*Offset=*/true> Item =
+              IDBuilder::createItem<Dims, true>(Range, ID, Offset);
+
+          if (StoreLocation) {
+            store_id(&ID);
+            store_item(&Item);
+          }
+          runKernelWithArg<const sycl::id<Dims> &>(MKernel, ID);
+        });
   }
 
   template <class ArgT = KernelArgType>
@@ -253,7 +345,7 @@ class HostKernel : public HostKernelBase {
         store_id(&ID);
         store_item(&ItemWithOffset);
       }
-      MKernel(Item);
+      runKernelWithArg<sycl::item<Dims, /*Offset=*/false>>(MKernel, Item);
     });
   }
 
@@ -276,18 +368,18 @@ class HostKernel : public HostKernelBase {
       UpperBound[I] = Range[I] + Offset[I];
     }
 
-    detail::NDLoop<Dims>::iterate(/*LowerBound=*/Offset, Stride, UpperBound,
-                                  [&](const sycl::id<Dims> &ID) {
-                                    sycl::item<Dims, /*Offset=*/true> Item =
-                                        IDBuilder::createItem<Dims, true>(
-                                            Range, ID, Offset);
-
-                                    if (StoreLocation) {
-                                      store_id(&ID);
-                                      store_item(&Item);
-                                    }
-                                    MKernel(Item);
-                                  });
+    detail::NDLoop<Dims>::iterate(
+        /*LowerBound=*/Offset, Stride, UpperBound,
+        [&](const sycl::id<Dims> &ID) {
+          sycl::item<Dims, /*Offset=*/true> Item =
+              IDBuilder::createItem<Dims, true>(Range, ID, Offset);
+
+          if (StoreLocation) {
+            store_id(&ID);
+            store_item(&Item);
+          }
+          runKernelWithArg<sycl::item<Dims, /*Offset=*/true>>(MKernel, Item);
+        });
   }
 
   template <class ArgT = KernelArgType>
@@ -336,7 +428,7 @@ class HostKernel : public HostKernelBase {
           auto g = NDItem.get_group();
           store_group(&g);
         }
-        MKernel(NDItem);
+        runKernelWithArg<const sycl::nd_item<Dims>>(MKernel, NDItem);
       });
     });
   }
@@ -364,7 +456,7 @@ class HostKernel : public HostKernelBase {
     detail::NDLoop<Dims>::iterate(NGroups, [&](const id<Dims> &GroupID) {
       sycl::group<Dims> Group =
           IDBuilder::createGroup<Dims>(GlobalSize, LocalSize, NGroups, GroupID);
-      MKernel(Group);
+      runKernelWithArg<sycl::group<Dims>>(MKernel, Group);
     });
   }
 
diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp
index 535c7069c2974..7fbe0d1dffe76 100644
--- a/sycl/include/CL/sycl/detail/kernel_desc.hpp
+++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp
@@ -26,7 +26,8 @@ enum class kernel_param_kind_t {
   kind_accessor = 0,
   kind_std_layout = 1, // standard layout object parameters
   kind_sampler = 2,
-  kind_pointer = 3
+  kind_pointer = 3,
+  kind_specialization_constants_buffer = 4,
 };
 
 // describes a kernel parameter
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index c16b151b64b7d..2523e917890e1 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -22,6 +22,7 @@
 #include <CL/sycl/item.hpp>
 #include <CL/sycl/kernel.hpp>
 #include <CL/sycl/kernel_bundle.hpp>
+#include <CL/sycl/kernel_handler.hpp>
 #include <CL/sycl/nd_item.hpp>
 #include <CL/sycl/nd_range.hpp>
 #include <CL/sycl/property_list.hpp>
@@ -116,27 +117,6 @@ template <typename Type> struct get_kernel_wrapper_name_t {
       typename get_kernel_name_t<detail::auto_name, Type>::name>;
 };
 
-template <typename, typename T> struct check_fn_signature {
-  static_assert(std::integral_constant<T, false>::value,
-                "Second template parameter is required to be of function type");
-};
-
-template <typename F, typename RetT, typename... Args>
-struct check_fn_signature<F, RetT(Args...)> {
-private:
-  template <typename T>
-  static constexpr auto check(T *) -> typename std::is_same<
-      decltype(std::declval<T>().operator()(std::declval<Args>()...)),
-      RetT>::type;
-
-  template <typename> static constexpr std::false_type check(...);
-
-  using type = decltype(check<F>(0));
-
-public:
-  static constexpr bool value = type::value;
-};
-
 __SYCL_EXPORT device getDeviceFromHandler(handler &);
 
 #if __SYCL_ID_QUERIES_FIT_IN_INT__
@@ -533,6 +513,13 @@ class __SYCL_EXPORT handler {
   template <typename KernelName, typename KernelType, int Dims,
             typename LambdaArgType>
   void StoreLambda(KernelType KernelFunc) {
+    if (detail::isKernelLambdaCallableWithKernelHandler<KernelType,
+                                                        LambdaArgType>() &&
+        MIsHost) {
+      throw cl::sycl::feature_not_supported(
+          "kernel_handler is not yet supported by host device.",
+          PI_INVALID_OPERATION);
+    }
     MHostKernel.reset(
         new detail::HostKernel<KernelType, LambdaArgType, Dims, KernelName>(
             KernelFunc));
@@ -822,17 +809,14 @@ class __SYCL_EXPORT handler {
       if (getenv("SYCL_PARALLEL_FOR_RANGE_ROUNDING_TRACE") != nullptr)
         std::cout << "parallel_for range adjusted from " << NumWorkItems[0]
                   << " to " << NewValX << std::endl;
-      auto Wrapper = [=](TransformedArgType Arg) {
-        if (Arg[0] >= NumWorkItems[0])
-          return;
-        Arg.set_allowed_range(NumWorkItems);
-        KernelFunc(Arg);
-      };
+
+      auto Wrapper = getRangeRoundedKernelLambda<TransformedArgType, Dims>(
+          KernelFunc, NumWorkItems);
 
       range<Dims> AdjustedRange = NumWorkItems;
       AdjustedRange.set_range_dim0(NewValX);
 #ifdef __SYCL_DEVICE_ONLY__
-      kernel_parallel_for<NameWT, TransformedArgType>(Wrapper);
+      kernel_parallel_for_wrapper<NameWT, TransformedArgType>(Wrapper);
 #else
       detail::checkValueRange<Dims>(AdjustedRange);
       MNDRDesc.set(std::move(AdjustedRange));
@@ -847,7 +831,7 @@ class __SYCL_EXPORT handler {
     {
 #ifdef __SYCL_DEVICE_ONLY__
       (void)NumWorkItems;
-      kernel_parallel_for<NameT, TransformedArgType>(KernelFunc);
+      kernel_parallel_for_wrapper<NameT, TransformedArgType>(KernelFunc);
 #else
       detail::checkValueRange<Dims>(NumWorkItems);
       MNDRDesc.set(std::move(NumWorkItems));
@@ -891,6 +875,18 @@ class __SYCL_EXPORT handler {
     KernelFunc();
   }
 
+  // NOTE: the name of this function - "kernel_single_task" - is used by the
+  // Front End to determine kernel invocation kind.
+  template <typename KernelName, typename KernelType>
+  __attribute__((sycl_kernel)) void
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_single_task(KernelType KernelFunc, kernel_handler KH) {
+#else
+  kernel_single_task(const KernelType &KernelFunc, kernel_handler KH) {
+#endif
+    KernelFunc(KH);
+  }
+
   // NOTE: the name of these functions - "kernel_parallel_for" - are used by the
   // Front End to determine kernel invocation kind.
   template <typename KernelName, typename ElementType, typename KernelType>
@@ -903,6 +899,18 @@ class __SYCL_EXPORT handler {
     KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
   }
 
+  // NOTE: the name of these functions - "kernel_parallel_for" - are used by the
+  // Front End to determine kernel invocation kind.
+  template <typename KernelName, typename ElementType, typename KernelType>
+  __attribute__((sycl_kernel)) void
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_parallel_for(KernelType KernelFunc, kernel_handler KH) {
+#else
+  kernel_parallel_for(const KernelType &KernelFunc, kernel_handler KH) {
+#endif
+    KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()), KH);
+  }
+
   // NOTE: the name of this function - "kernel_parallel_for_work_group" - is
   // used by the Front End to determine kernel invocation kind.
   template <typename KernelName, typename ElementType, typename KernelType>
@@ -915,6 +923,103 @@ class __SYCL_EXPORT handler {
     KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
   }
 
+  // NOTE: the name of this function - "kernel_parallel_for_work_group" - is
+  // used by the Front End to determine kernel invocation kind.
+  template <typename KernelName, typename ElementType, typename KernelType>
+  __attribute__((sycl_kernel)) void
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_parallel_for_work_group(KernelType KernelFunc, kernel_handler KH) {
+#else
+  kernel_parallel_for_work_group(const KernelType &KernelFunc,
+                                 kernel_handler KH) {
+#endif
+    KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()), KH);
+  }
+
+  // Wrappers for kernel_*** functions above with and without support of
+  // additional kernel_handler argument.
+
+  // NOTE: to support kernel_handler argument in kernel lambdas, only
+  // kernel_***_wrapper functions must be called in this code
+
+  // Wrappers for kernel_single_task(...)
+
+  template <typename KernelName, typename KernelType>
+  std::enable_if_t<
+      detail::isKernelLambdaCallableWithKernelHandler<KernelType>(), void>
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_single_task_wrapper(KernelType KernelFunc) {
+#else
+  kernel_single_task_wrapper(const KernelType &KernelFunc) {
+#endif
+    kernel_handler KH;
+    kernel_single_task<KernelName>(KernelFunc, KH);
+  }
+
+  template <typename KernelName, typename KernelType>
+  std::enable_if_t<
+      !detail::isKernelLambdaCallableWithKernelHandler<KernelType>(), void>
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_single_task_wrapper(KernelType KernelFunc) {
+#else
+  kernel_single_task_wrapper(const KernelType &KernelFunc) {
+#endif
+    kernel_single_task<KernelName>(KernelFunc);
+  }
+
+  // Wrappers for kernel_parallel_for(...)
+
+  template <typename KernelName, typename ElementType, typename KernelType>
+  std::enable_if_t<detail::isKernelLambdaCallableWithKernelHandler<
+                       KernelType, ElementType>(),
+                   void>
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_parallel_for_wrapper(KernelType KernelFunc) {
+#else
+  kernel_parallel_for_wrapper(const KernelType &KernelFunc) {
+#endif
+    kernel_handler KH;
+    kernel_parallel_for<KernelName, ElementType>(KernelFunc, KH);
+  }
+
+  template <typename KernelName, typename ElementType, typename KernelType>
+  std::enable_if_t<!detail::isKernelLambdaCallableWithKernelHandler<
+                       KernelType, ElementType>(),
+                   void>
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_parallel_for_wrapper(KernelType KernelFunc) {
+#else
+  kernel_parallel_for_wrapper(const KernelType &KernelFunc) {
+#endif
+    kernel_parallel_for<KernelName, ElementType>(KernelFunc);
+  }
+
+  // Wrappers for kernel_parallel_for_work_group(...)
+
+  template <typename KernelName, typename ElementType, typename KernelType>
+  std::enable_if_t<detail::isKernelLambdaCallableWithKernelHandler<
+                       KernelType, ElementType>(),
+                   void>
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_parallel_for_work_group_wrapper(KernelType KernelFunc) {
+#else
+  kernel_parallel_for_work_group_wrapper(const KernelType &KernelFunc) {
+#endif
+    kernel_handler KH;
+    kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc, KH);
+  }
+
+  template <typename KernelName, typename ElementType, typename KernelType>
+  std::enable_if_t<!detail::isKernelLambdaCallableWithKernelHandler<
+                       KernelType, ElementType>(),
+                   void>
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_parallel_for_work_group_wrapper(KernelType KernelFunc) {
+#else
+  kernel_parallel_for_work_group_wrapper(const KernelType &KernelFunc) {
+#endif
+    kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc);
+  }
 #endif
 
   std::shared_ptr<detail::kernel_bundle_impl>
@@ -1061,7 +1166,7 @@ class __SYCL_EXPORT handler {
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
-    kernel_single_task<NameT>(KernelFunc);
+    kernel_single_task_wrapper<NameT>(KernelFunc);
 #else
     // No need to check if range is out of INT_MAX limits as it's compile-time
     // known constant.
@@ -1175,7 +1280,7 @@ class __SYCL_EXPORT handler {
 #ifdef __SYCL_DEVICE_ONLY__
     (void)NumWorkItems;
     (void)WorkItemOffset;
-    kernel_parallel_for<NameT, LambdaArgType>(KernelFunc);
+    kernel_parallel_for_wrapper<NameT, LambdaArgType>(KernelFunc);
 #else
     detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
     MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
@@ -1207,7 +1312,7 @@ class __SYCL_EXPORT handler {
         sycl::detail::lambda_arg_type<KernelType, nd_item<Dims>>;
 #ifdef __SYCL_DEVICE_ONLY__
     (void)ExecutionRange;
-    kernel_parallel_for<NameT, LambdaArgType>(KernelFunc);
+    kernel_parallel_for_wrapper<NameT, LambdaArgType>(KernelFunc);
 #else
     detail::checkValueRange<Dims>(ExecutionRange);
     MNDRDesc.set(std::move(ExecutionRange));
@@ -1440,7 +1545,7 @@ class __SYCL_EXPORT handler {
         sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
 #ifdef __SYCL_DEVICE_ONLY__
     (void)NumWorkGroups;
-    kernel_parallel_for_work_group<NameT, LambdaArgType>(KernelFunc);
+    kernel_parallel_for_work_group_wrapper<NameT, LambdaArgType>(KernelFunc);
 #else
     detail::checkValueRange<Dims>(NumWorkGroups);
     MNDRDesc.setNumWorkGroups(NumWorkGroups);
@@ -1474,7 +1579,7 @@ class __SYCL_EXPORT handler {
 #ifdef __SYCL_DEVICE_ONLY__
     (void)NumWorkGroups;
     (void)WorkGroupSize;
-    kernel_parallel_for_work_group<NameT, LambdaArgType>(KernelFunc);
+    kernel_parallel_for_work_group_wrapper<NameT, LambdaArgType>(KernelFunc);
 #else
     nd_range<Dims> ExecRange =
         nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
@@ -1610,7 +1715,7 @@ class __SYCL_EXPORT handler {
 #ifdef __SYCL_DEVICE_ONLY__
     (void)Kernel;
     (void)NumWorkItems;
-    kernel_parallel_for<NameT, LambdaArgType>(KernelFunc);
+    kernel_parallel_for_wrapper<NameT, LambdaArgType>(KernelFunc);
 #else
     detail::checkValueRange<Dims>(NumWorkItems);
     MNDRDesc.set(std::move(NumWorkItems));
@@ -1646,7 +1751,7 @@ class __SYCL_EXPORT handler {
     (void)Kernel;
     (void)NumWorkItems;
     (void)WorkItemOffset;
-    kernel_parallel_for<NameT, LambdaArgType>(KernelFunc);
+    kernel_parallel_for_wrapper<NameT, LambdaArgType>(KernelFunc);
 #else
     detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
     MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
@@ -1682,7 +1787,7 @@ class __SYCL_EXPORT handler {
 #ifdef __SYCL_DEVICE_ONLY__
     (void)Kernel;
     (void)NDRange;
-    kernel_parallel_for<NameT, LambdaArgType>(KernelFunc);
+    kernel_parallel_for_wrapper<NameT, LambdaArgType>(KernelFunc);
 #else
     detail::checkValueRange<Dims>(NDRange);
     MNDRDesc.set(std::move(NDRange));
@@ -1722,7 +1827,7 @@ class __SYCL_EXPORT handler {
 #ifdef __SYCL_DEVICE_ONLY__
     (void)Kernel;
     (void)NumWorkGroups;
-    kernel_parallel_for_work_group<NameT, LambdaArgType>(KernelFunc);
+    kernel_parallel_for_work_group_wrapper<NameT, LambdaArgType>(KernelFunc);
 #else
     detail::checkValueRange<Dims>(NumWorkGroups);
     MNDRDesc.setNumWorkGroups(NumWorkGroups);
@@ -1761,7 +1866,7 @@ class __SYCL_EXPORT handler {
     (void)Kernel;
     (void)NumWorkGroups;
     (void)WorkGroupSize;
-    kernel_parallel_for_work_group<NameT, LambdaArgType>(KernelFunc);
+    kernel_parallel_for_work_group_wrapper<NameT, LambdaArgType>(KernelFunc);
 #else
     nd_range<Dims> ExecRange =
         nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
@@ -2170,6 +2275,34 @@ class __SYCL_EXPORT handler {
                                            access::target);
 
   friend class ::MockHandler;
+
+  template <
+      typename TransformedArgType, int Dims, typename KernelType,
+      typename std::enable_if_t<detail::isKernelLambdaCallableWithKernelHandler<
+          KernelType, TransformedArgType>()> * = nullptr>
+  auto getRangeRoundedKernelLambda(KernelType KernelFunc,
+                                   range<Dims> NumWorkItems) {
+    return [=](TransformedArgType Arg, kernel_handler KH) {
+      if (Arg[0] >= NumWorkItems[0])
+        return;
+      Arg.set_allowed_range(NumWorkItems);
+      KernelFunc(Arg, KH);
+    };
+  }
+
+  template <typename TransformedArgType, int Dims, typename KernelType,
+            typename std::enable_if_t<
+                !detail::isKernelLambdaCallableWithKernelHandler<
+                    KernelType, TransformedArgType>()> * = nullptr>
+  auto getRangeRoundedKernelLambda(KernelType KernelFunc,
+                                   range<Dims> NumWorkItems) {
+    return [=](TransformedArgType Arg) {
+      if (Arg[0] >= NumWorkItems[0])
+        return;
+      Arg.set_allowed_range(NumWorkItems);
+      KernelFunc(Arg);
+    };
+  }
 };
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/kernel_handler.hpp b/sycl/include/CL/sycl/kernel_handler.hpp
new file mode 100644
index 0000000000000..0a3f77381b364
--- /dev/null
+++ b/sycl/include/CL/sycl/kernel_handler.hpp
@@ -0,0 +1,73 @@
+//==------ kernel_handler.hpp -- SYCL standard header file -----*- C++ -*---==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+// This guard is needed because the libsycl.so can compiled with C++ <=14
+// while the code requires C++17. This code is not supposed to be used by the
+// libsycl.so so it should not be a problem.
+#if __cplusplus > 201402L
+template <auto &S> struct specialization_id_name_generator {};
+#endif
+
+} // namespace detail
+
+/// Reading the value of a specialization constant
+///
+/// \ingroup sycl_api
+class kernel_handler {
+public:
+#if __cplusplus > 201402L
+  template <auto &S>
+  typename std::remove_reference_t<decltype(S)> get_specialization_constant() {
+#ifdef __SYCL_DEVICE_ONLY__
+    return getSpecializationConstantOnDevice<S>();
+#else
+    // TODO: add support of host device
+    throw cl::sycl::feature_not_supported(
+        "kernel_handler::get_specialization_constant() is not yet supported by "
+        "host device.",
+        PI_INVALID_OPERATION);
+#endif // __SYCL_DEVICE_ONLY__
+  }
+#endif // __cplusplus > 201402L
+
+private:
+  void __init_specialization_constants_buffer(
+      char *SpecializationConstantsBuffer = nullptr) {
+    MSpecializationConstantsBuffer = SpecializationConstantsBuffer;
+  }
+
+#ifdef __SYCL_DEVICE_ONLY__
+  template <auto &S, typename T = std::remove_reference_t<decltype(S)>,
+            std::enable_if_t<std::is_fundamental_v<T>> * = nullptr>
+  T getSpecializationConstantOnDevice() {
+    const char *SymbolicID = __builtin_unique_stable_name(
+        detail::specialization_id_name_generator<S>);
+    return __sycl_getScalar2020SpecConstantValue<T>(
+        SymbolicID, &S, MSpecializationConstantsBuffer);
+  }
+  template <auto &S, typename T = std::remove_reference_t<decltype(S)>,
+            std::enable_if_t<std::is_compound_v<T>> * = nullptr>
+  T getSpecializationConstantOnDevice() {
+    const char *SymbolicID = __builtin_unique_stable_name(
+        detail::specialization_id_name_generator<S>);
+    return __sycl_getComposite2020SpecConstantValue<T>(
+        SymbolicID, &S, MSpecializationConstantsBuffer);
+  }
+#endif // __SYCL_DEVICE_ONLY__
+
+  char *MSpecializationConstantsBuffer = nullptr;
+};
+
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/specialization_id.hpp b/sycl/include/CL/sycl/specialization_id.hpp
new file mode 100644
index 0000000000000..cc4f25d8ebad6
--- /dev/null
+++ b/sycl/include/CL/sycl/specialization_id.hpp
@@ -0,0 +1,35 @@
+//==---- specialization_id.hpp -- SYCL standard header file ----*- C++ -*---==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+
+/// Declaring a specialization constant
+///
+/// \ingroup sycl_api
+template <typename T> class specialization_id {
+public:
+  using value_type = T;
+
+  template <class... Args>
+  explicit constexpr specialization_id(Args &&... args)
+      : MDefaultValue(std::forward<Args>(args)...) {}
+
+  specialization_id(const specialization_id &rhs) = delete;
+  specialization_id(specialization_id &&rhs) = delete;
+  specialization_id &operator=(const specialization_id &rhs) = delete;
+  specialization_id &operator=(specialization_id &&rhs) = delete;
+
+private:
+  T MDefaultValue;
+};
+
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp
index 99b6f13a4c823..1072c23fbdafd 100644
--- a/sycl/source/detail/scheduler/commands.cpp
+++ b/sycl/source/detail/scheduler/commands.cpp
@@ -1696,6 +1696,12 @@ pi_result ExecCGCommand::SetKernelParamsAndLaunch(
                                                        Arg.MSize, Arg.MPtr);
       break;
     }
+    case kernel_param_kind_t::kind_specialization_constants_buffer: {
+      throw cl::sycl::feature_not_supported(
+          "SYCL2020 specialization constants are not yet fully supported",
+          PI_INVALID_OPERATION);
+      break;
+    }
     }
     ++NextTrueIndex;
   }
diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp
index 68ad4f9094d1e..9ba4887ca935f 100644
--- a/sycl/source/handler.cpp
+++ b/sycl/source/handler.cpp
@@ -335,6 +335,12 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
                        Index + IndexShift);
     break;
   }
+  case kernel_param_kind_t::kind_specialization_constants_buffer: {
+    throw cl::sycl::feature_not_supported(
+        "SYCL2020 specialization constants are not yet fully supported",
+        PI_INVALID_OPERATION);
+    break;
+  }
   }
 }
 
diff --git a/sycl/test/function-pointers/fp-as-kernel-arg.cpp b/sycl/test/function-pointers/fp-as-kernel-arg.cpp
index 5973c2b419bc8..65b6d0c6be988 100644
--- a/sycl/test/function-pointers/fp-as-kernel-arg.cpp
+++ b/sycl/test/function-pointers/fp-as-kernel-arg.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out
+// RUN: %clangxx -Xclang -fsycl-allow-func-ptr -fsycl %s -o %t.out
 // RUN: %RUN_ON_HOST %t.out
 // FIXME: This test should use runtime early exit once correct check for
 // corresponding extension is implemented
diff --git a/sycl/test/function-pointers/pass-fp-through-buffer.cpp b/sycl/test/function-pointers/pass-fp-through-buffer.cpp
index 60ca4cd5f4c71..aef94716ecd85 100644
--- a/sycl/test/function-pointers/pass-fp-through-buffer.cpp
+++ b/sycl/test/function-pointers/pass-fp-through-buffer.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out
+// RUN: %clangxx -Xclang -fsycl-allow-func-ptr -fsycl %s -o %t.out
 // RUN: %RUN_ON_HOST %t.out
 // FIXME: This test should use runtime early exit once correct check for
 // corresponding extension is implemented
diff --git a/sycl/test/on-device/basic_tests/specialization_constants/kernel_lambda_with_kernel_handler_arg.cpp b/sycl/test/on-device/basic_tests/specialization_constants/kernel_lambda_with_kernel_handler_arg.cpp
new file mode 100644
index 0000000000000..7adf90fe48b7a
--- /dev/null
+++ b/sycl/test/on-device/basic_tests/specialization_constants/kernel_lambda_with_kernel_handler_arg.cpp
@@ -0,0 +1,112 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %t.out
+
+// This test checks all possible scenarios of running single_task, parallel_for
+// and parallel_for_work_group to verify that this code compiles and runs
+// correctly with user's lambda with and without sycl::kernel_handler argument
+
+// TODO: enable cuda support when non-native spec constants started to be
+// supported
+// UNSUPPORTED: cuda
+
+#include <CL/sycl.hpp>
+
+int main() {
+  sycl::queue q;
+
+  // single_task w/o kernel_handler arg
+  q.submit([&](sycl::handler &cgh) {
+    cgh.single_task<class KernelSingleTaskWOKernelHandler>([=]() {});
+  });
+
+  // single_task with kernel_handler arg
+  q.submit([&](sycl::handler &cgh) {
+    cgh.single_task<class KernelSingleTaskWithKernelHandler>(
+        [=](sycl::kernel_handler kh) {});
+  });
+
+  // parallel_for with id and w/o kernel_handler arg
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for<class KernelParallelForIdWithoutKernelHandler>(
+        sycl::range<1>(1), [](sycl::id<1> i) {});
+  });
+
+  // parallel_for with id and kernel_handler args
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for<class KernelParallelForIdWithKernelHandler>(
+        sycl::range<1>(1), [](sycl::id<1> i, sycl::kernel_handler kh) {});
+  });
+
+  // parallel_for with item and w/o kernel_handler arg
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for<class KernelParallelForItemWithoutKernelHandler>(
+        sycl::range<3>(3, 3, 3), [](sycl::item<3> it) {});
+  });
+
+  // parallel_for with item and kernel_handler args
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for<class KernelParallelForItemWithKernelHandler>(
+        sycl::range<3>(3, 3, 3),
+        [](sycl::item<3> it, sycl::kernel_handler kh) {});
+  });
+
+  // parallel_for with nd_item and w/o kernel_handler arg
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for<class KernelParallelForNDItemWithoutKernelHandler>(
+        sycl::nd_range<3>(sycl::range<3>(4, 4, 4), sycl::range<3>(2, 2, 2)),
+        [=](sycl::nd_item<3> item) {});
+  });
+
+  // parallel_for with nd_item and kernel_handler args
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for<class KernelParallelForNDItemWithKernelHandler>(
+        sycl::nd_range<3>(sycl::range<3>(4, 4, 4), sycl::range<3>(2, 2, 2)),
+        [=](sycl::nd_item<3> item, sycl::kernel_handler kh) {});
+  });
+
+  // parallel_for with generic lambda w/o kernel_handler arg
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for<class KernelParallelForGenericLambdaWithoutKernelHandler>(
+        sycl::range<3>(3, 3, 3), [](auto it) {});
+  });
+
+  // parallel_for with generic lambda with kernel_handler arg
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for<class KernelParallelForGenericLambdaWithKernelHandler>(
+        sycl::range<3>(3, 3, 3), [](auto it, sycl::kernel_handler kh) {});
+  });
+
+  // parallel_for with integral type arg and w/o kernel_handler arg
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for<class KernelParallelForIntWithoutKernelHandler>(
+        sycl::range<1>(1), [](int index) {});
+  });
+
+  // parallel_for with integral type and kernel_handler args
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for<class KernelParallelForIntWithKernelHandler>(
+        sycl::range<1>(1), [](int index, sycl::kernel_handler kh) {});
+  });
+
+  // parallel_for_work_group without kernel_handler arg
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for_work_group<
+        class KernelParallelForWorkGroupWithKernelHandler>(
+        sycl::range<3>(2, 2, 2), sycl::range<3>(2, 2, 2),
+        [=](sycl::group<3> myGroup) {
+          myGroup.parallel_for_work_item([&](sycl::h_item<3> myItem) {});
+          myGroup.parallel_for_work_item([&](sycl::h_item<3> myItem) {});
+        });
+  });
+
+  // parallel_for_work_group with kernel_handler arg
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for_work_group<
+        class KernelParallelForWorkGroupWithoutKernelHandler>(
+        sycl::range<3>(2, 2, 2), sycl::range<3>(2, 2, 2),
+        [=](sycl::group<3> myGroup, sycl::kernel_handler kh) {
+          myGroup.parallel_for_work_item([&](sycl::h_item<3> myItem) {});
+          myGroup.parallel_for_work_item([&](sycl::h_item<3> myItem) {});
+        });
+  });
+}
diff --git a/sycl/test/regression/sycl-include-gnu14.cpp b/sycl/test/regression/sycl-include-gnu17.cpp
similarity index 74%
rename from sycl/test/regression/sycl-include-gnu14.cpp
rename to sycl/test/regression/sycl-include-gnu17.cpp
index 0147acf143de5..f56d262c1f0bd 100644
--- a/sycl/test/regression/sycl-include-gnu14.cpp
+++ b/sycl/test/regression/sycl-include-gnu17.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -std=gnu++14 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -std=gnu++17 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %RUN_ON_HOST %t.out
 
 // UNSUPPORTED: system-windows