From 6f79af1c84adac613fc927497595bfea81ed9cb7 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Mon, 21 Dec 2020 17:44:46 -0800
Subject: [PATCH 01/13] initial commit for time trials.

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 sycl/include/CL/sycl/ONEAPI/reduction.hpp |  1 +
 sycl/include/CL/sycl/detail/cg_types.hpp  | 61 ++++++++++++++++++-----
 sycl/include/CL/sycl/handler.hpp          | 14 +++---
 sycl/include/CL/sycl/kernel.hpp           | 16 +++++-
 4 files changed, 73 insertions(+), 19 deletions(-)

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..965362984a10c 100644
--- a/sycl/include/CL/sycl/detail/cg_types.hpp
+++ b/sycl/include/CL/sycl/detail/cg_types.hpp
@@ -11,6 +11,7 @@
 #include <CL/sycl/detail/host_profiling_info.hpp>
 #include <CL/sycl/detail/kernel_desc.hpp>
 #include <CL/sycl/group.hpp>
+//#include <CL/sycl/handler.hpp>
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/interop_handle.hpp>
 #include <CL/sycl/interop_handler.hpp>
@@ -158,7 +159,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 +204,12 @@ 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) {
+    //CP
+    using NameT = typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    std::string KName = typeid(NameT *).name();
+    using KI = detail::KernelInfo<KernelName>;
+    constexpr bool StoreLocation = KI::callsThisItem(); //TO callsThisID or ThisItem?
+
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
     sycl::id<Dims> Offset;
     for (int I = 0; I < Dims; ++I) {
@@ -213,8 +220,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);
+      //CP
+      if(StoreLocation){
+        store_id(&ID);  // <--
+        store_item(&Item);
+      }
       MKernel(ID);
     });
   }
@@ -223,6 +233,12 @@ class HostKernel : public HostKernelBase {
   typename detail::enable_if_t<
       std::is_same<ArgT, item<Dims, /*Offset=*/false>>::value>
   runOnHost(const NDRDescT &NDRDesc) {
+    //CP
+    using NameT = typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    std::string KName = typeid(NameT *).name();
+    using KI = detail::KernelInfo<KernelName>;
+    constexpr bool StoreLocation = KI::callsThisItem(); //TO callsThisID or ThisItem?
+
     sycl::id<Dims> ID;
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
     for (int I = 0; I < Dims; ++I)
@@ -232,8 +248,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);
+      //CP
+      if(StoreLocation){
+        store_id(&ID);
+        store_item(&ItemWithOffset);
+      }
       MKernel(Item);
     });
   }
@@ -242,6 +261,12 @@ class HostKernel : public HostKernelBase {
   typename detail::enable_if_t<
       std::is_same<ArgT, item<Dims, /*Offset=*/true>>::value>
   runOnHost(const NDRDescT &NDRDesc) {
+    //CP
+    using NameT = typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    std::string KName = typeid(NameT *).name();
+    using KI = detail::KernelInfo<KernelName>;
+    constexpr bool StoreLocation = KI::callsThisItem(); //TO callsThisID or ThisItem?
+
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
     sycl::id<Dims> Offset;
     for (int I = 0; I < Dims; ++I) {
@@ -253,8 +278,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);
+      //CP
+      if(StoreLocation){
+        store_id(&OffsetID);
+        store_item(&Item);
+      }
       MKernel(Item);
     });
   }
@@ -262,6 +290,12 @@ 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) {
+    //CP
+    using NameT = typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    std::string KName = typeid(NameT *).name();
+    using KI = detail::KernelInfo<KernelName>;
+    constexpr bool StoreLocation = KI::callsThisItem(); //TO callsThisID or ThisItem?
+
     sycl::range<Dims> GroupSize(InitializedVal<Dims, range>::template get<0>());
     for (int I = 0; I < Dims; ++I) {
       if (NDRDesc.LocalSize[I] == 0 ||
@@ -294,11 +328,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);
+        //CP    
+        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/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 117d700596072..de4c437d4df50 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -78,9 +78,10 @@ template <typename T, int Dimensions, typename AllocatorT, typename Enable>
 class buffer;
 namespace detail {
 
+//CP
 /// This class is the default KernelName template parameter type for kernel
 /// invocation APIs such as single_task.
-class auto_name {};
+//class auto_name {};
 
 class kernel_impl;
 class queue_impl;
@@ -111,12 +112,13 @@ SuggestedArgType argument_helper(...);
 template <typename F, typename SuggestedArgType>
 using lambda_arg_type = decltype(argument_helper<F, SuggestedArgType>(0));
 
+//CP
 /// 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;
-};
+// 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> {
@@ -508,7 +510,7 @@ 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 +1057,7 @@ 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..ff91ae178e485 100644
--- a/sycl/include/CL/sycl/kernel.hpp
+++ b/sycl/include/CL/sycl/kernel.hpp
@@ -22,7 +22,21 @@ class program;
 class context;
 namespace detail {
 class kernel_impl;
-}
+
+
+//CP
+/// 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;
+};
+
+} //detail namespace
 
 /// Provides an abstraction of a SYCL kernel.
 ///

From a30131c44eb7bf0f05fd33f20abfa3e68ce2b45c Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 22 Dec 2020 17:43:28 -0800
Subject: [PATCH 02/13] expanding callsThisItem functionality to include all
 four free functions and checking for any usage to optimize host kernel tasks.

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/include/clang/Sema/Sema.h             | 21 +++++++++-
 clang/lib/Sema/SemaSYCL.cpp                 | 44 +++++++++++++++++++--
 sycl/include/CL/sycl/detail/cg_types.hpp    | 20 ++++------
 sycl/include/CL/sycl/detail/kernel_desc.hpp |  2 +
 sycl/include/CL/sycl/handler.hpp            | 11 ------
 sycl/include/CL/sycl/kernel.hpp             |  1 -
 6 files changed, 71 insertions(+), 28 deletions(-)

diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 4528044881604..49ee0e76f5f2c 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -346,8 +346,17 @@ class SYCLIntegrationHeader {
   /// Registers a specialization constant to emit info for it into the header.
   void addSpecConstant(StringRef IDName, QualType IDType);
 
+  // CP
   /// Notes that this_item is called within the kernel.
+  // void setCallsThisItem(bool B);
+
+  // CP
+  /// 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 +375,15 @@ class SYCLIntegrationHeader {
     KernelParamDesc() = default;
   };
 
+  // there are four free function 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.
@@ -386,7 +404,8 @@ class SYCLIntegrationHeader {
     SmallVector<KernelParamDesc, 8> Params;
 
     // Whether kernel calls this_item()
-    bool CallsThisItem;
+    // bool CallsThisItem;
+    KernelCallsSYCLFreeFunction FreeFunctionCalls;
 
     KernelDesc() = default;
   };
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 9ad45432935e1..adc00bb8d8190 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 kernels");
+  K->FreeFunctionCalls.CallsThisId = B;
+}
+
 void SYCLIntegrationHeader::setCallsThisItem(bool B) {
   KernelDesc *K = getCurKernelDesc();
   assert(K && "no kernels");
-  K->CallsThisItem = B;
+  K->FreeFunctionCalls.CallsThisItem = B;
+}
+
+void SYCLIntegrationHeader::setCallsThisNDItem(bool B) {
+  KernelDesc *K = getCurKernelDesc();
+  assert(K && "no kernels");
+  K->FreeFunctionCalls.CallsThisNDItem = B;
+}
+
+void SYCLIntegrationHeader::setCallsThisGroup(bool B) {
+  KernelDesc *K = getCurKernelDesc();
+  assert(K && "no kernels");
+  K->FreeFunctionCalls.CallsThisGroup = B;
 }
 
 SYCLIntegrationHeader::SYCLIntegrationHeader(DiagnosticsEngine &_Diag,
diff --git a/sycl/include/CL/sycl/detail/cg_types.hpp b/sycl/include/CL/sycl/detail/cg_types.hpp
index 965362984a10c..32863af67c72f 100644
--- a/sycl/include/CL/sycl/detail/cg_types.hpp
+++ b/sycl/include/CL/sycl/detail/cg_types.hpp
@@ -204,11 +204,10 @@ 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) {
-    //CP
     using NameT = typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
-    constexpr bool StoreLocation = KI::callsThisItem(); //TO callsThisID or ThisItem?
+    constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
 
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
     sycl::id<Dims> Offset;
@@ -220,7 +219,7 @@ 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);
-      //CP
+          
       if(StoreLocation){
         store_id(&ID);  // <--
         store_item(&Item);
@@ -233,11 +232,10 @@ class HostKernel : public HostKernelBase {
   typename detail::enable_if_t<
       std::is_same<ArgT, item<Dims, /*Offset=*/false>>::value>
   runOnHost(const NDRDescT &NDRDesc) {
-    //CP
     using NameT = typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
-    constexpr bool StoreLocation = KI::callsThisItem(); //TO callsThisID or ThisItem?
+    constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
 
     sycl::id<Dims> ID;
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
@@ -248,7 +246,7 @@ class HostKernel : public HostKernelBase {
       sycl::item<Dims, /*Offset=*/false> Item =
           IDBuilder::createItem<Dims, false>(Range, ID);
       sycl::item<Dims, /*Offset=*/true> ItemWithOffset = Item;
-      //CP
+
       if(StoreLocation){
         store_id(&ID);
         store_item(&ItemWithOffset);
@@ -261,11 +259,10 @@ class HostKernel : public HostKernelBase {
   typename detail::enable_if_t<
       std::is_same<ArgT, item<Dims, /*Offset=*/true>>::value>
   runOnHost(const NDRDescT &NDRDesc) {
-    //CP
     using NameT = typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
-    constexpr bool StoreLocation = KI::callsThisItem(); //TO callsThisID or ThisItem?
+    constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
 
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
     sycl::id<Dims> Offset;
@@ -278,7 +275,7 @@ class HostKernel : public HostKernelBase {
       sycl::id<Dims> OffsetID = ID + Offset;
       sycl::item<Dims, /*Offset=*/true> Item =
           IDBuilder::createItem<Dims, true>(Range, OffsetID, Offset);
-      //CP
+
       if(StoreLocation){
         store_id(&OffsetID);
         store_item(&Item);
@@ -290,11 +287,10 @@ 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) {
-    //CP
     using NameT = typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
-    constexpr bool StoreLocation = KI::callsThisItem(); //TO callsThisID or ThisItem?
+    constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
 
     sycl::range<Dims> GroupSize(InitializedVal<Dims, range>::template get<0>());
     for (int I = 0; I < Dims; ++I) {
@@ -328,7 +324,7 @@ class HostKernel : public HostKernelBase {
             IDBuilder::createItem<Dims, false>(LocalSize, LocalID);
         const sycl::nd_item<Dims> NDItem =
             IDBuilder::createNDItem<Dims>(GlobalItem, LocalItem, Group);
-        //CP    
+
         if(StoreLocation){
           store_id(&GlobalID);
           store_item(&GlobalItem);
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 de4c437d4df50..d026c356ad630 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 {
 
-//CP
-/// 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;
@@ -112,13 +108,6 @@ SuggestedArgType argument_helper(...);
 template <typename F, typename SuggestedArgType>
 using lambda_arg_type = decltype(argument_helper<F, SuggestedArgType>(0));
 
-//CP
-/// 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> {
diff --git a/sycl/include/CL/sycl/kernel.hpp b/sycl/include/CL/sycl/kernel.hpp
index ff91ae178e485..ad1c1cd9c4a9c 100644
--- a/sycl/include/CL/sycl/kernel.hpp
+++ b/sycl/include/CL/sycl/kernel.hpp
@@ -24,7 +24,6 @@ namespace detail {
 class kernel_impl;
 
 
-//CP
 /// This class is the default KernelName template parameter type for kernel
 /// invocation APIs such as single_task.
 class auto_name {};

From 15d5bf4e427e827947a4f39637e6bdd8fb06b2ea Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 29 Dec 2020 11:12:54 -0800
Subject: [PATCH 03/13] updated tests

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/test/CodeGenSYCL/Inputs/sycl.hpp        |  3 ++
 .../CodeGenSYCL/parallel_for_this_item.cpp    | 33 +++++++++++++++++--
 2 files changed, 34 insertions(+), 2 deletions(-)

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/parallel_for_this_item.cpp b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
index 422a1bad33373..362a08bd444f3 100755
--- a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
+++ b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -sycl-std=2020 -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
 // RUN: FileCheck -input-file=%t.h %s
 
 // This test checks that compiler generates correct kernel description
@@ -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;

From 3419fb653bedd4aca747600f7682ba8a4d6ac361 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 29 Dec 2020 11:49:11 -0800
Subject: [PATCH 04/13] overlooked comment

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/include/clang/Sema/Sema.h | 5 -----
 1 file changed, 5 deletions(-)

diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 49ee0e76f5f2c..b71745061c953 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -346,11 +346,6 @@ class SYCLIntegrationHeader {
   /// Registers a specialization constant to emit info for it into the header.
   void addSpecConstant(StringRef IDName, QualType IDType);
 
-  // CP
-  /// Notes that this_item is called within the kernel.
-  // void setCallsThisItem(bool B);
-
-  // CP
   /// Note which free functions (this_id, this_item, etc) are called within the
   /// kernel
   void setCallsThisId(bool B);

From 9d2d8439ae20bfe8d451f5fc03148353e43fadc9 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 29 Dec 2020 11:57:29 -0800
Subject: [PATCH 05/13] clang-format continues to declare its devotion. This is
 my tender reply.

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 sycl/include/CL/sycl/detail/cg_types.hpp | 24 ++++++++++++++----------
 sycl/include/CL/sycl/handler.hpp         | 12 ++++++------
 sycl/include/CL/sycl/kernel.hpp          |  3 +--
 3 files changed, 21 insertions(+), 18 deletions(-)

diff --git a/sycl/include/CL/sycl/detail/cg_types.hpp b/sycl/include/CL/sycl/detail/cg_types.hpp
index 32863af67c72f..c79d056f6dcf3 100644
--- a/sycl/include/CL/sycl/detail/cg_types.hpp
+++ b/sycl/include/CL/sycl/detail/cg_types.hpp
@@ -204,7 +204,8 @@ 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 NameT = typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    using NameT =
+        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
     constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
@@ -219,9 +220,9 @@ 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);
-          
-      if(StoreLocation){
-        store_id(&ID);  // <--
+
+      if (StoreLocation) {
+        store_id(&ID);
         store_item(&Item);
       }
       MKernel(ID);
@@ -232,7 +233,8 @@ class HostKernel : public HostKernelBase {
   typename detail::enable_if_t<
       std::is_same<ArgT, item<Dims, /*Offset=*/false>>::value>
   runOnHost(const NDRDescT &NDRDesc) {
-    using NameT = typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    using NameT =
+        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
     constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
@@ -247,7 +249,7 @@ class HostKernel : public HostKernelBase {
           IDBuilder::createItem<Dims, false>(Range, ID);
       sycl::item<Dims, /*Offset=*/true> ItemWithOffset = Item;
 
-      if(StoreLocation){
+      if (StoreLocation) {
         store_id(&ID);
         store_item(&ItemWithOffset);
       }
@@ -259,7 +261,8 @@ class HostKernel : public HostKernelBase {
   typename detail::enable_if_t<
       std::is_same<ArgT, item<Dims, /*Offset=*/true>>::value>
   runOnHost(const NDRDescT &NDRDesc) {
-    using NameT = typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    using NameT =
+        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
     constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
@@ -276,7 +279,7 @@ class HostKernel : public HostKernelBase {
       sycl::item<Dims, /*Offset=*/true> Item =
           IDBuilder::createItem<Dims, true>(Range, OffsetID, Offset);
 
-      if(StoreLocation){
+      if (StoreLocation) {
         store_id(&OffsetID);
         store_item(&Item);
       }
@@ -287,7 +290,8 @@ 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 NameT = typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    using NameT =
+        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
     constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
@@ -325,7 +329,7 @@ class HostKernel : public HostKernelBase {
         const sycl::nd_item<Dims> NDItem =
             IDBuilder::createNDItem<Dims>(GlobalItem, LocalItem, Group);
 
-        if(StoreLocation){
+        if (StoreLocation) {
           store_id(&GlobalID);
           store_item(&GlobalItem);
           store_nd_item(&NDItem);
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index d026c356ad630..a1e14501cee92 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -78,7 +78,6 @@ template <typename T, int Dimensions, typename AllocatorT, typename Enable>
 class buffer;
 namespace detail {
 
-
 class kernel_impl;
 class queue_impl;
 class stream_impl;
@@ -108,7 +107,6 @@ SuggestedArgType argument_helper(...);
 template <typename F, typename SuggestedArgType>
 using lambda_arg_type = decltype(argument_helper<F, SuggestedArgType>(0));
 
-
 /// Specialization for the case when \c Name is undefined.
 template <typename Type> struct get_kernel_name_t<detail::auto_name, Type> {
   using name = Type;
@@ -400,7 +398,7 @@ class __SYCL_EXPORT handler {
   // Recursively calls itself until arguments pack is fully processed.
   // The version for regular(standard layout) argument.
   template <typename T, typename... Ts>
-  void setArgsHelper(int ArgIndex, T &&Arg, Ts &&... Args) {
+  void setArgsHelper(int ArgIndex, T &&Arg, Ts &&...Args) {
     set_arg(ArgIndex, std::move(Arg));
     setArgsHelper(++ArgIndex, std::move(Args)...);
   }
@@ -499,7 +497,8 @@ class __SYCL_EXPORT handler {
             typename LambdaArgType>
   void StoreLambda(KernelType KernelFunc) {
     MHostKernel.reset(
-        new detail::HostKernel<KernelType, LambdaArgType, Dims, KernelName>(KernelFunc));
+        new detail::HostKernel<KernelType, LambdaArgType, Dims, KernelName>(
+            KernelFunc));
 
     using KI = sycl::detail::KernelInfo<KernelName>;
     // Empty name indicates that the compilation happens without integration
@@ -976,7 +975,7 @@ class __SYCL_EXPORT handler {
   /// Registers pack of arguments(Args) with indexes starting from 0.
   ///
   /// \param Args are argument values to be set.
-  template <typename... Ts> void set_args(Ts &&... Args) {
+  template <typename... Ts> void set_args(Ts &&...Args) {
     setArgsHelper(0, std::move(Args)...);
   }
 
@@ -1046,7 +1045,8 @@ class __SYCL_EXPORT handler {
     MNDRDesc.set(range<1>{1});
 
     MArgs = std::move(MAssociatedAccesors);
-    MHostKernel.reset(new detail::HostKernel<FuncT, void, 1, void>(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 ad1c1cd9c4a9c..a02ab2a4ef5d7 100644
--- a/sycl/include/CL/sycl/kernel.hpp
+++ b/sycl/include/CL/sycl/kernel.hpp
@@ -23,7 +23,6 @@ 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 {};
@@ -35,7 +34,7 @@ template <typename Name, typename Type> struct get_kernel_name_t {
   using name = Name;
 };
 
-} //detail namespace
+} // namespace detail
 
 /// Provides an abstraction of a SYCL kernel.
 ///

From a05c6676e3800629e2119ddd6469d57e520973aa Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 29 Dec 2020 12:04:31 -0800
Subject: [PATCH 06/13] clang-format appreciates little my penitent reply

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 sycl/include/CL/sycl/handler.hpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index a1e14501cee92..cd48679e96644 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -398,7 +398,7 @@ class __SYCL_EXPORT handler {
   // Recursively calls itself until arguments pack is fully processed.
   // The version for regular(standard layout) argument.
   template <typename T, typename... Ts>
-  void setArgsHelper(int ArgIndex, T &&Arg, Ts &&...Args) {
+  void setArgsHelper(int ArgIndex, T &&Arg, Ts &&... Args) {
     set_arg(ArgIndex, std::move(Arg));
     setArgsHelper(++ArgIndex, std::move(Args)...);
   }
@@ -975,7 +975,7 @@ class __SYCL_EXPORT handler {
   /// Registers pack of arguments(Args) with indexes starting from 0.
   ///
   /// \param Args are argument values to be set.
-  template <typename... Ts> void set_args(Ts &&...Args) {
+  template <typename... Ts> void set_args(Ts &&... Args) {
     setArgsHelper(0, std::move(Args)...);
   }
 

From 8526e5e367b6611ed13e4461dbfc69ae26e8c7e8 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 29 Dec 2020 13:36:03 -0800
Subject: [PATCH 07/13] restore original test invocation line

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/test/CodeGenSYCL/parallel_for_this_item.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
index 362a08bd444f3..dc484b33bfbb4 100755
--- a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
+++ b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -sycl-std=2020 -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
 // RUN: FileCheck -input-file=%t.h %s
 
 // This test checks that compiler generates correct kernel description

From da71fe91625f88e4a6b2f2cc0b763ec480608e5d Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 29 Dec 2020 14:32:30 -0800
Subject: [PATCH 08/13] updated expectation line no. due to change in proxy
 sycl.hpp

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/test/CodeGenSYCL/kernel-by-reference.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

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) {

From 0ecd023a27f9764e051fac5731314f3f35a95c31 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 30 Dec 2020 09:47:57 -0800
Subject: [PATCH 09/13] changes requested by reviewers

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/include/clang/Sema/Sema.h          |  2 +-
 clang/lib/Sema/SemaSYCL.cpp              | 10 +++++-----
 sycl/include/CL/sycl/detail/cg_types.hpp |  1 -
 3 files changed, 6 insertions(+), 7 deletions(-)

diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index b71745061c953..f2553ee1acef1 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -370,7 +370,7 @@ class SYCLIntegrationHeader {
     KernelParamDesc() = default;
   };
 
-  // there are four free function the kernel may call (this_id, this_item,
+  // there are four free functions the kernel may call (this_id, this_item,
   // this_nd_item, this_group)
   struct KernelCallsSYCLFreeFunction {
     bool CallsThisId;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index adc00bb8d8190..94b82ab3cc6cf 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -2721,7 +2721,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
       if (!Visited.insert(FD).second)
         continue; // We've already seen this Decl
 
-      // Check whether this call is to free functions (  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);
@@ -4001,25 +4001,25 @@ void SYCLIntegrationHeader::addSpecConstant(StringRef IDName, QualType IDType) {
 
 void SYCLIntegrationHeader::setCallsThisId(bool B) {
   KernelDesc *K = getCurKernelDesc();
-  assert(K && "no kernels");
+  assert(K && "no kernel");
   K->FreeFunctionCalls.CallsThisId = B;
 }
 
 void SYCLIntegrationHeader::setCallsThisItem(bool B) {
   KernelDesc *K = getCurKernelDesc();
-  assert(K && "no kernels");
+  assert(K && "no kernel");
   K->FreeFunctionCalls.CallsThisItem = B;
 }
 
 void SYCLIntegrationHeader::setCallsThisNDItem(bool B) {
   KernelDesc *K = getCurKernelDesc();
-  assert(K && "no kernels");
+  assert(K && "no kernel");
   K->FreeFunctionCalls.CallsThisNDItem = B;
 }
 
 void SYCLIntegrationHeader::setCallsThisGroup(bool B) {
   KernelDesc *K = getCurKernelDesc();
-  assert(K && "no kernels");
+  assert(K && "no kernel");
   K->FreeFunctionCalls.CallsThisGroup = B;
 }
 
diff --git a/sycl/include/CL/sycl/detail/cg_types.hpp b/sycl/include/CL/sycl/detail/cg_types.hpp
index c79d056f6dcf3..3210f59abe9e8 100644
--- a/sycl/include/CL/sycl/detail/cg_types.hpp
+++ b/sycl/include/CL/sycl/detail/cg_types.hpp
@@ -11,7 +11,6 @@
 #include <CL/sycl/detail/host_profiling_info.hpp>
 #include <CL/sycl/detail/kernel_desc.hpp>
 #include <CL/sycl/group.hpp>
-//#include <CL/sycl/handler.hpp>
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/interop_handle.hpp>
 #include <CL/sycl/interop_handler.hpp>

From 28412524ec1be3aecc73c5c596942c35862a2fb2 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 30 Dec 2020 11:39:42 -0800
Subject: [PATCH 10/13] reviewer requests

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/include/clang/Sema/Sema.h | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index f2553ee1acef1..c017080381471 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -398,8 +398,7 @@ 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;

From c43160fb0bbc9def47c447577d21ebc6db03dfa6 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 30 Dec 2020 12:25:07 -0800
Subject: [PATCH 11/13] clang-format <3  4-evah

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/include/clang/Sema/Sema.h | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index c017080381471..64f110713d030 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -398,7 +398,8 @@ class SYCLIntegrationHeader {
     /// Descriptor of kernel actual parameters.
     SmallVector<KernelParamDesc, 8> Params;
 
-    // Whether kernel calls any of the SYCL free functions (this_item(), this_id(), etc)
+    // Whether kernel calls any of the SYCL free functions (this_item(),
+    // this_id(), etc)
     KernelCallsSYCLFreeFunction FreeFunctionCalls;
 
     KernelDesc() = default;

From 80453fda5a28d85007829fc23baef22e011d7b85 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 5 Jan 2021 11:16:25 -0800
Subject: [PATCH 12/13] changes requested in review

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 sycl/include/CL/sycl/detail/cg_types.hpp | 12 ------------
 1 file changed, 12 deletions(-)

diff --git a/sycl/include/CL/sycl/detail/cg_types.hpp b/sycl/include/CL/sycl/detail/cg_types.hpp
index 3210f59abe9e8..11826f7e6989c 100644
--- a/sycl/include/CL/sycl/detail/cg_types.hpp
+++ b/sycl/include/CL/sycl/detail/cg_types.hpp
@@ -203,9 +203,6 @@ 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 NameT =
-        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
-    std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
     constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
 
@@ -232,9 +229,6 @@ class HostKernel : public HostKernelBase {
   typename detail::enable_if_t<
       std::is_same<ArgT, item<Dims, /*Offset=*/false>>::value>
   runOnHost(const NDRDescT &NDRDesc) {
-    using NameT =
-        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
-    std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
     constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
 
@@ -260,9 +254,6 @@ class HostKernel : public HostKernelBase {
   typename detail::enable_if_t<
       std::is_same<ArgT, item<Dims, /*Offset=*/true>>::value>
   runOnHost(const NDRDescT &NDRDesc) {
-    using NameT =
-        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
-    std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
     constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
 
@@ -289,9 +280,6 @@ 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 NameT =
-        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
-    std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
     constexpr bool StoreLocation = KI::callsAnyThisFreeFunction();
 

From f05da87db940ea2dc41b1978e0c90f179a29ba84 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 5 Jan 2021 11:34:39 -0800
Subject: [PATCH 13/13] more requested changes. put get_kernel_name_t decl in
 one place

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 sycl/include/CL/sycl/handler.hpp | 5 -----
 sycl/include/CL/sycl/kernel.hpp  | 5 +++++
 2 files changed, 5 insertions(+), 5 deletions(-)

diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index cd48679e96644..6b13ba100fd92 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -107,11 +107,6 @@ SuggestedArgType argument_helper(...);
 template <typename F, typename SuggestedArgType>
 using lambda_arg_type = decltype(argument_helper<F, SuggestedArgType>(0));
 
-/// 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;
 
diff --git a/sycl/include/CL/sycl/kernel.hpp b/sycl/include/CL/sycl/kernel.hpp
index a02ab2a4ef5d7..3370112892c6c 100644
--- a/sycl/include/CL/sycl/kernel.hpp
+++ b/sycl/include/CL/sycl/kernel.hpp
@@ -34,6 +34,11 @@ 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.