diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 4528044881604..64f110713d030 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -346,8 +346,12 @@ class SYCLIntegrationHeader {
   /// Registers a specialization constant to emit info for it into the header.
   void addSpecConstant(StringRef IDName, QualType IDType);
 
-  /// Notes that this_item is called within the kernel.
+  /// Note which free functions (this_id, this_item, etc) are called within the
+  /// kernel
+  void setCallsThisId(bool B);
   void setCallsThisItem(bool B);
+  void setCallsThisNDItem(bool B);
+  void setCallsThisGroup(bool B);
 
 private:
   // Kernel actual parameter descriptor.
@@ -366,6 +370,15 @@ class SYCLIntegrationHeader {
     KernelParamDesc() = default;
   };
 
+  // there are four free functions the kernel may call (this_id, this_item,
+  // this_nd_item, this_group)
+  struct KernelCallsSYCLFreeFunction {
+    bool CallsThisId;
+    bool CallsThisItem;
+    bool CallsThisNDItem;
+    bool CallsThisGroup;
+  };
+
   // Kernel invocation descriptor
   struct KernelDesc {
     /// Kernel name.
@@ -385,8 +398,9 @@ class SYCLIntegrationHeader {
     /// Descriptor of kernel actual parameters.
     SmallVector<KernelParamDesc, 8> Params;
 
-    // Whether kernel calls this_item()
-    bool CallsThisItem;
+    // Whether kernel calls any of the SYCL free functions (this_item(),
+    // this_id(), etc)
+    KernelCallsSYCLFreeFunction FreeFunctionCalls;
 
     KernelDesc() = default;
   };
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 9ad45432935e1..94b82ab3cc6cf 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -2721,11 +2721,24 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
       if (!Visited.insert(FD).second)
         continue; // We've already seen this Decl
 
-      // Check whether this call is to sycl::this_item().
+      // Check whether this call is to free functions (sycl::this_item(),
+      // this_id, etc.).
+      if (Util::isSyclFunction(FD, "this_id")) {
+        Header.setCallsThisId(true);
+        return;
+      }
       if (Util::isSyclFunction(FD, "this_item")) {
         Header.setCallsThisItem(true);
         return;
       }
+      if (Util::isSyclFunction(FD, "this_nd_item")) {
+        Header.setCallsThisNDItem(true);
+        return;
+      }
+      if (Util::isSyclFunction(FD, "this_group")) {
+        Header.setCallsThisGroup(true);
+        return;
+      }
 
       CallGraphNode *N = SYCLCG.getNode(FD);
       if (!N)
@@ -3920,7 +3933,14 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
       << "; }\n";
     O << "  __SYCL_DLL_LOCAL\n";
     O << "  static constexpr bool callsThisItem() { return ";
-    O << K.CallsThisItem << "; }\n";
+    O << K.FreeFunctionCalls.CallsThisItem << "; }\n";
+    O << "  __SYCL_DLL_LOCAL\n";
+    O << "  static constexpr bool callsAnyThisFreeFunction() { return ";
+    O << (K.FreeFunctionCalls.CallsThisId ||
+          K.FreeFunctionCalls.CallsThisItem ||
+          K.FreeFunctionCalls.CallsThisNDItem ||
+          K.FreeFunctionCalls.CallsThisGroup)
+      << "; }\n";
     O << "};\n";
     CurStart += N;
   }
@@ -3979,10 +3999,28 @@ void SYCLIntegrationHeader::addSpecConstant(StringRef IDName, QualType IDType) {
   SpecConsts.emplace_back(std::make_pair(IDType, IDName.str()));
 }
 
+void SYCLIntegrationHeader::setCallsThisId(bool B) {
+  KernelDesc *K = getCurKernelDesc();
+  assert(K && "no kernel");
+  K->FreeFunctionCalls.CallsThisId = B;
+}
+
 void SYCLIntegrationHeader::setCallsThisItem(bool B) {
   KernelDesc *K = getCurKernelDesc();
-  assert(K && "no kernels");
-  K->CallsThisItem = B;
+  assert(K && "no kernel");
+  K->FreeFunctionCalls.CallsThisItem = B;
+}
+
+void SYCLIntegrationHeader::setCallsThisNDItem(bool B) {
+  KernelDesc *K = getCurKernelDesc();
+  assert(K && "no kernel");
+  K->FreeFunctionCalls.CallsThisNDItem = B;
+}
+
+void SYCLIntegrationHeader::setCallsThisGroup(bool B) {
+  KernelDesc *K = getCurKernelDesc();
+  assert(K && "no kernel");
+  K->FreeFunctionCalls.CallsThisGroup = B;
 }
 
 SYCLIntegrationHeader::SYCLIntegrationHeader(DiagnosticsEngine &_Diag,
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index 0f71db428018a..5462a2d13d1f6 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -130,6 +130,9 @@ template <int dim> struct item {
 template <int Dims> item<Dims>
 this_item() { return item<Dims>{}; }
 
+template <int Dims> id<Dims>
+this_id() { return id<Dims>{}; }
+
 template <int dim>
 struct range {
   template <typename... T>
diff --git a/clang/test/CodeGenSYCL/kernel-by-reference.cpp b/clang/test/CodeGenSYCL/kernel-by-reference.cpp
index f5bbac0e75730..23a09c43413a9 100644
--- a/clang/test/CodeGenSYCL/kernel-by-reference.cpp
+++ b/clang/test/CodeGenSYCL/kernel-by-reference.cpp
@@ -15,7 +15,7 @@ int simple_add(int i) {
 int main() {
   queue q;
 #if defined(SYCL2020)
-  // expected-warning@Inputs/sycl.hpp:298 {{Passing kernel functions by value is deprecated in SYCL 2020}}
+  // expected-warning@Inputs/sycl.hpp:301 {{Passing kernel functions by value is deprecated in SYCL 2020}}
   // expected-note@+3 {{in instantiation of function template specialization}}
 #endif
   q.submit([&](handler &h) {
@@ -23,7 +23,7 @@ int main() {
   });
 
 #if defined(SYCL2017)
-  // expected-warning@Inputs/sycl.hpp:293 {{Passing of kernel functions by reference is a SYCL 2020 extension}}
+  // expected-warning@Inputs/sycl.hpp:296 {{Passing of kernel functions by reference is a SYCL 2020 extension}}
   // expected-note@+3 {{in instantiation of function template specialization}}
 #endif
   q.submit([&](handler &h) {
diff --git a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
index 422a1bad33373..dc484b33bfbb4 100755
--- a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
+++ b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
@@ -13,7 +13,8 @@
 // CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3GNU",
 // CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3EMU",
 // CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3OWL",
-// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3RAT"
+// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3RAT",
+// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3FOX"
 // CHECK-NEXT: };
 
 // CHECK:template <> struct KernelInfo<class GNU> {
@@ -29,6 +30,8 @@
 // CHECK-NEXT:  static constexpr bool isESIMD() { return 0; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
 // CHECK-NEXT:  static constexpr bool callsThisItem() { return 0; }
+// CHECK-NEXT:  __SYCL_DLL_LOCAL
+// CHECK-NEXT:  static constexpr bool callsAnyThisFreeFunction() { return 0; }
 // CHECK-NEXT:};
 // CHECK-NEXT:template <> struct KernelInfo<class EMU> {
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
@@ -43,6 +46,8 @@
 // CHECK-NEXT:  static constexpr bool isESIMD() { return 0; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
 // CHECK-NEXT:    static constexpr bool callsThisItem() { return 1; }
+// CHECK-NEXT:  __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr bool callsAnyThisFreeFunction() { return 1; }
 // CHECK-NEXT:};
 // CHECK-NEXT:template <> struct KernelInfo<class OWL> {
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
@@ -57,6 +62,8 @@
 // CHECK-NEXT:  static constexpr bool isESIMD() { return 0; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
 // CHECK-NEXT:    static constexpr bool callsThisItem() { return 0; }
+// CHECK-NEXT:  __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr bool callsAnyThisFreeFunction() { return 0; }
 // CHECK-NEXT:};
 // CHECK-NEXT:template <> struct KernelInfo<class RAT> {
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
@@ -71,6 +78,24 @@
 // CHECK-NEXT:  static constexpr bool isESIMD() { return 0; }
 // CHECK-NEXT: __SYCL_DLL_LOCAL
 // CHECK-NEXT:    static constexpr bool callsThisItem() { return 1; }
+// CHECK-NEXT: __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr bool callsAnyThisFreeFunction() { return 1; }
+// CHECK-NEXT:};
+// CHECK-NEXT:template <> struct KernelInfo<class FOX> {
+// CHECK-NEXT:  __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3FOX"; }
+// CHECK-NEXT:  __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr unsigned getNumParams() { return 0; }
+// CHECK-NEXT:  __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr const kernel_param_desc_t& getParamDesc(unsigned i) {
+// CHECK-NEXT:    return kernel_signatures[i+0];
+// CHECK-NEXT: }
+// CHECK-NEXT:  __SYCL_DLL_LOCAL
+// CHECK-NEXT:  static constexpr bool isESIMD() { return 0; }
+// CHECK-NEXT: __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr bool callsThisItem() { return 0; }
+// CHECK-NEXT: __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr bool callsAnyThisFreeFunction() { return 1; }
 // CHECK-NEXT:};
 
 #include "sycl.hpp"
@@ -108,6 +133,10 @@ int main() {
 
     // This kernel calls sycl::this_item
     cgh.parallel_for<class RAT>(range<1>(1), [=](id<1> I) { f(); });
+
+    // This kernel does not call sycl::this_item, but does call this_id
+    cgh.parallel_for<class FOX>(range<1>(1),
+                                [=](id<1> I) { this_id<1>(); });
   });
 
   return 0;
diff --git a/sycl/include/CL/sycl/ONEAPI/reduction.hpp b/sycl/include/CL/sycl/ONEAPI/reduction.hpp
index a0c415373b447..f15a811c21873 100644
--- a/sycl/include/CL/sycl/ONEAPI/reduction.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/reduction.hpp
@@ -12,6 +12,7 @@
 #include <CL/sycl/ONEAPI/group_algorithm.hpp>
 #include <CL/sycl/accessor.hpp>
 #include <CL/sycl/handler.hpp>
+#include <CL/sycl/kernel.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
diff --git a/sycl/include/CL/sycl/detail/cg_types.hpp b/sycl/include/CL/sycl/detail/cg_types.hpp
index 7e14e0aed035f..11826f7e6989c 100644
--- a/sycl/include/CL/sycl/detail/cg_types.hpp
+++ b/sycl/include/CL/sycl/detail/cg_types.hpp
@@ -158,7 +158,7 @@ class HostTask {
 };
 
 // Class which stores specific lambda object.
-template <class KernelType, class KernelArgType, int Dims>
+template <class KernelType, class KernelArgType, int Dims, typename KernelName>
 class HostKernel : public HostKernelBase {
   using IDBuilder = sycl::detail::Builder;
   KernelType MKernel;
@@ -203,6 +203,9 @@ class HostKernel : public HostKernelBase {
   template <class ArgT = KernelArgType>
   typename detail::enable_if_t<std::is_same<ArgT, sycl::id<Dims>>::value>
   runOnHost(const NDRDescT &NDRDesc) {
+    using KI = detail::KernelInfo<KernelName>;
+    constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
+
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
     sycl::id<Dims> Offset;
     for (int I = 0; I < Dims; ++I) {
@@ -213,8 +216,11 @@ class HostKernel : public HostKernelBase {
     detail::NDLoop<Dims>::iterate(Range, [&](const sycl::id<Dims> &ID) {
       sycl::item<Dims, /*Offset=*/true> Item =
           IDBuilder::createItem<Dims, true>(Range, ID, Offset);
-      store_id(&ID);
-      store_item(&Item);
+
+      if (StoreLocation) {
+        store_id(&ID);
+        store_item(&Item);
+      }
       MKernel(ID);
     });
   }
@@ -223,6 +229,9 @@ class HostKernel : public HostKernelBase {
   typename detail::enable_if_t<
       std::is_same<ArgT, item<Dims, /*Offset=*/false>>::value>
   runOnHost(const NDRDescT &NDRDesc) {
+    using KI = detail::KernelInfo<KernelName>;
+    constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
+
     sycl::id<Dims> ID;
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
     for (int I = 0; I < Dims; ++I)
@@ -232,8 +241,11 @@ class HostKernel : public HostKernelBase {
       sycl::item<Dims, /*Offset=*/false> Item =
           IDBuilder::createItem<Dims, false>(Range, ID);
       sycl::item<Dims, /*Offset=*/true> ItemWithOffset = Item;
-      store_id(&ID);
-      store_item(&ItemWithOffset);
+
+      if (StoreLocation) {
+        store_id(&ID);
+        store_item(&ItemWithOffset);
+      }
       MKernel(Item);
     });
   }
@@ -242,6 +254,9 @@ class HostKernel : public HostKernelBase {
   typename detail::enable_if_t<
       std::is_same<ArgT, item<Dims, /*Offset=*/true>>::value>
   runOnHost(const NDRDescT &NDRDesc) {
+    using KI = detail::KernelInfo<KernelName>;
+    constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
+
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
     sycl::id<Dims> Offset;
     for (int I = 0; I < Dims; ++I) {
@@ -253,8 +268,11 @@ class HostKernel : public HostKernelBase {
       sycl::id<Dims> OffsetID = ID + Offset;
       sycl::item<Dims, /*Offset=*/true> Item =
           IDBuilder::createItem<Dims, true>(Range, OffsetID, Offset);
-      store_id(&OffsetID);
-      store_item(&Item);
+
+      if (StoreLocation) {
+        store_id(&OffsetID);
+        store_item(&Item);
+      }
       MKernel(Item);
     });
   }
@@ -262,6 +280,9 @@ class HostKernel : public HostKernelBase {
   template <class ArgT = KernelArgType>
   typename detail::enable_if_t<std::is_same<ArgT, nd_item<Dims>>::value>
   runOnHost(const NDRDescT &NDRDesc) {
+    using KI = detail::KernelInfo<KernelName>;
+    constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
+
     sycl::range<Dims> GroupSize(InitializedVal<Dims, range>::template get<0>());
     for (int I = 0; I < Dims; ++I) {
       if (NDRDesc.LocalSize[I] == 0 ||
@@ -294,11 +315,14 @@ class HostKernel : public HostKernelBase {
             IDBuilder::createItem<Dims, false>(LocalSize, LocalID);
         const sycl::nd_item<Dims> NDItem =
             IDBuilder::createNDItem<Dims>(GlobalItem, LocalItem, Group);
-        store_id(&GlobalID);
-        store_item(&GlobalItem);
-        store_nd_item(&NDItem);
-        auto g = NDItem.get_group();
-        store_group(&g);
+
+        if (StoreLocation) {
+          store_id(&GlobalID);
+          store_item(&GlobalItem);
+          store_nd_item(&NDItem);
+          auto g = NDItem.get_group();
+          store_group(&g);
+        }
         MKernel(NDItem);
       });
     });
diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp
index f3bf02b1b1492..c2c201fb4ecf5 100644
--- a/sycl/include/CL/sycl/detail/kernel_desc.hpp
+++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp
@@ -58,6 +58,7 @@ template <class KernelNameType> struct KernelInfo {
   static constexpr const char *getName() { return ""; }
   static constexpr bool isESIMD() { return 0; }
   static constexpr bool callsThisItem() { return false; }
+  static constexpr bool callsAnyThisFreeFunction() { return false; }
 };
 #else
 template <char...> struct KernelInfoData {
@@ -69,6 +70,7 @@ template <char...> struct KernelInfoData {
   static constexpr const char *getName() { return ""; }
   static constexpr bool isESIMD() { return 0; }
   static constexpr bool callsThisItem() { return false; }
+  static constexpr bool callsAnyThisFreeFunction() { return false; }
 };
 
 // C++14 like index_sequence and make_index_sequence
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 117d700596072..6b13ba100fd92 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -78,10 +78,6 @@ template <typename T, int Dimensions, typename AllocatorT, typename Enable>
 class buffer;
 namespace detail {
 
-/// This class is the default KernelName template parameter type for kernel
-/// invocation APIs such as single_task.
-class auto_name {};
-
 class kernel_impl;
 class queue_impl;
 class stream_impl;
@@ -111,18 +107,6 @@ SuggestedArgType argument_helper(...);
 template <typename F, typename SuggestedArgType>
 using lambda_arg_type = decltype(argument_helper<F, SuggestedArgType>(0));
 
-/// Helper struct to get a kernel name type based on given \c Name and \c Type
-/// types: if \c Name is undefined (is a \c auto_name) then \c Type becomes
-/// the \c Name.
-template <typename Name, typename Type> struct get_kernel_name_t {
-  using name = Name;
-};
-
-/// Specialization for the case when \c Name is undefined.
-template <typename Type> struct get_kernel_name_t<detail::auto_name, Type> {
-  using name = Type;
-};
-
 // Used when parallel_for range is rounded-up.
 template <typename Type> class __pf_kernel_wrapper;
 
@@ -508,7 +492,8 @@ class __SYCL_EXPORT handler {
             typename LambdaArgType>
   void StoreLambda(KernelType KernelFunc) {
     MHostKernel.reset(
-        new detail::HostKernel<KernelType, LambdaArgType, Dims>(KernelFunc));
+        new detail::HostKernel<KernelType, LambdaArgType, Dims, KernelName>(
+            KernelFunc));
 
     using KI = sycl::detail::KernelInfo<KernelName>;
     // Empty name indicates that the compilation happens without integration
@@ -1055,7 +1040,8 @@ class __SYCL_EXPORT handler {
     MNDRDesc.set(range<1>{1});
 
     MArgs = std::move(MAssociatedAccesors);
-    MHostKernel.reset(new detail::HostKernel<FuncT, void, 1>(std::move(Func)));
+    MHostKernel.reset(
+        new detail::HostKernel<FuncT, void, 1, void>(std::move(Func)));
     MCGType = detail::CG::RUN_ON_HOST_INTEL;
   }
 
diff --git a/sycl/include/CL/sycl/kernel.hpp b/sycl/include/CL/sycl/kernel.hpp
index c71a4cf510025..3370112892c6c 100644
--- a/sycl/include/CL/sycl/kernel.hpp
+++ b/sycl/include/CL/sycl/kernel.hpp
@@ -22,7 +22,24 @@ class program;
 class context;
 namespace detail {
 class kernel_impl;
-}
+
+/// This class is the default KernelName template parameter type for kernel
+/// invocation APIs such as single_task.
+class auto_name {};
+
+/// Helper struct to get a kernel name type based on given \c Name and \c Type
+/// types: if \c Name is undefined (is a \c auto_name) then \c Type becomes
+/// the \c Name.
+template <typename Name, typename Type> struct get_kernel_name_t {
+  using name = Name;
+};
+
+/// Specialization for the case when \c Name is undefined.
+template <typename Type> struct get_kernel_name_t<detail::auto_name, Type> {
+  using name = Type;
+};
+
+} // namespace detail
 
 /// Provides an abstraction of a SYCL kernel.
 ///