diff --git a/clang/test/SemaSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp b/clang/test/SemaSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp
index c1812a0324868..6f797e0f93604 100644
--- a/clang/test/SemaSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp
+++ b/clang/test/SemaSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp
@@ -35,6 +35,17 @@ __SYCL_INLINE_NAMESPACE(cl) {
     int offset;
   };
 
+  template <bool Cond, typename TrueT, typename FalseT>
+  struct conditional {
+    using type = TrueT;
+  };
+  template <typename TrueT, typename FalseT>
+  struct conditional<false, TrueT, FalseT> {
+    using type = FalseT;
+  };
+
+  using int64_t = conditional<sizeof(long) == 8, long, long long>::type;
+
   template <class KernelNameType> struct KernelInfo {
     static constexpr unsigned getNumParams() { return 0; }
     static const kernel_param_desc_t &getParamDesc(int) {
@@ -43,6 +54,7 @@ __SYCL_INLINE_NAMESPACE(cl) {
     }
     static constexpr const char *getName() { return ""; }
     static constexpr bool isESIMD() { return 0; }
+    static constexpr int64_t getKernelSize() { return 0; }
   };
   } // namespace detail
   } // namespace sycl
diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp
index 3545631083a70..243d7b7f75e01 100644
--- a/sycl/include/CL/sycl/detail/kernel_desc.hpp
+++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp
@@ -13,6 +13,8 @@
 #include <CL/sycl/detail/defines_elementary.hpp>
 #include <CL/sycl/detail/export.hpp>
 
+#include <cstdint>
+
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
@@ -79,6 +81,7 @@ template <class KernelNameType> struct KernelInfo {
   static constexpr const char *getFunctionName() { return ""; }
   static constexpr unsigned getLineNumber() { return 0; }
   static constexpr unsigned getColumnNumber() { return 0; }
+  static constexpr int64_t getKernelSize() { return 0; }
 };
 #else
 template <char...> struct KernelInfoData {
@@ -93,6 +96,7 @@ template <char...> struct KernelInfoData {
   static constexpr const char *getFunctionName() { return ""; }
   static constexpr unsigned getLineNumber() { return 0; }
   static constexpr unsigned getColumnNumber() { return 0; }
+  static constexpr int64_t getKernelSize() { return 0; }
 };
 
 // C++14 like index_sequence and make_index_sequence
@@ -135,6 +139,9 @@ template <class KernelNameType> struct KernelInfo {
   static constexpr const char *getFunctionName() { return ""; }
   static constexpr unsigned getLineNumber() { return 0; }
   static constexpr unsigned getColumnNumber() { return 0; }
+  static constexpr int64_t getKernelSize() {
+    return SubKernelInfo::getKernelSize();
+  }
 };
 #endif //__SYCL_UNNAMED_LAMBDA__
 
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 7d2dbea3ad279..5bc0ae2c11acc 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -717,13 +717,25 @@ class __SYCL_EXPORT handler {
           "kernel_handler is not yet supported by host device.",
           PI_ERROR_INVALID_OPERATION);
     }
+
     KernelType *KernelPtr =
         ResetHostKernel<KernelType, LambdaArgType, Dims>(KernelFunc);
 
     using KI = sycl::detail::KernelInfo<KernelName>;
+    constexpr bool KernelHasName =
+        KI::getName() != nullptr && KI::getName()[0] != '\0';
+
+    // Some host compilers may have different captures from Clang. Currently
+    // there is no stable way of handling this when extracting the captures, so
+    // a static assert is made to fail for incompatible kernel lambdas.
+    static_assert(!KernelHasName || sizeof(KernelFunc) == KI::getKernelSize(),
+                  "Unexpected kernel lambda size. This can be caused by an "
+                  "external host compiler producing a lambda with an "
+                  "unexpected layout. This is a limitation of the compiler.");
+
     // Empty name indicates that the compilation happens without integration
     // header, so don't perform things that require it.
-    if (KI::getName() != nullptr && KI::getName()[0] != '\0') {
+    if (KernelHasName) {
       // TODO support ESIMD in no-integration-header case too.
       MArgs.clear();
       extractArgsAndReqsFromLambda(reinterpret_cast<char *>(KernelPtr),
diff --git a/sycl/test/basic_tests/kernel_size_mismatch.cpp b/sycl/test/basic_tests/kernel_size_mismatch.cpp
new file mode 100644
index 0000000000000..c8ea60d575c63
--- /dev/null
+++ b/sycl/test/basic_tests/kernel_size_mismatch.cpp
@@ -0,0 +1,19 @@
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning -o - %s
+
+// Tests for static assertion failure when kernel lambda mismatches between host
+// and device.
+
+#include <CL/sycl.hpp>
+
+int main() {
+  sycl::queue Q;
+  int A = 1;
+  Q.single_task([=]() {
+#ifdef __SYCL_DEVICE_ONLY__
+     (void)A;
+    // expected-no-diagnostics
+#else
+  // expected-error-re@CL/sycl/handler.hpp:* {{static_assert failed due to requirement '{{.*}}' "Unexpected kernel lambda size. This can be caused by an external host compiler producing a lambda with an unexpected layout. This is a limitation of the compiler."}}
+#endif
+   }).wait();
+}
diff --git a/sycl/unittests/SYCL2020/GetNativeOpenCL.cpp b/sycl/unittests/SYCL2020/GetNativeOpenCL.cpp
index 751d4bca1e7a7..bad193277d964 100644
--- a/sycl/unittests/SYCL2020/GetNativeOpenCL.cpp
+++ b/sycl/unittests/SYCL2020/GetNativeOpenCL.cpp
@@ -127,7 +127,8 @@ TEST(GetNative, GetNativeHandle) {
   sycl::buffer<int, 1> Buffer(&Data[0], sycl::range<1>(1));
   Queue.submit([&](sycl::handler &cgh) {
     auto Acc = Buffer.get_access<sycl::access::mode::read_write>(cgh);
-    cgh.single_task<TestKernel>([=]() { (void)Acc; });
+    constexpr size_t KS = sizeof(decltype(Acc));
+    cgh.single_task<TestKernel<KS>>([=]() { (void)Acc; });
   });
 
   get_native<backend::opencl>(Context);
diff --git a/sycl/unittests/SYCL2020/KernelBundle.cpp b/sycl/unittests/SYCL2020/KernelBundle.cpp
index b7d33c4d8bcec..c25e397770e3d 100644
--- a/sycl/unittests/SYCL2020/KernelBundle.cpp
+++ b/sycl/unittests/SYCL2020/KernelBundle.cpp
@@ -31,6 +31,7 @@ template <> struct KernelInfo<TestKernel> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 template <> struct KernelInfo<TestKernelExeOnly> {
@@ -43,6 +44,7 @@ template <> struct KernelInfo<TestKernelExeOnly> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 } // namespace detail
diff --git a/sycl/unittests/SYCL2020/KernelID.cpp b/sycl/unittests/SYCL2020/KernelID.cpp
index add7f9d1097f5..8c122692e5d72 100644
--- a/sycl/unittests/SYCL2020/KernelID.cpp
+++ b/sycl/unittests/SYCL2020/KernelID.cpp
@@ -32,6 +32,7 @@ template <> struct KernelInfo<TestKernel1> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 template <> struct KernelInfo<TestKernel2> {
@@ -44,6 +45,7 @@ template <> struct KernelInfo<TestKernel2> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 template <> struct KernelInfo<TestKernel3> {
@@ -56,6 +58,7 @@ template <> struct KernelInfo<TestKernel3> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 template <> struct KernelInfo<ServiceKernel1> {
@@ -70,6 +73,7 @@ template <> struct KernelInfo<ServiceKernel1> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 } // namespace detail
 } // namespace sycl
diff --git a/sycl/unittests/SYCL2020/SpecializationConstant.cpp b/sycl/unittests/SYCL2020/SpecializationConstant.cpp
index b64e53eadb694..3c511f28f27db 100644
--- a/sycl/unittests/SYCL2020/SpecializationConstant.cpp
+++ b/sycl/unittests/SYCL2020/SpecializationConstant.cpp
@@ -35,6 +35,7 @@ template <> struct KernelInfo<TestKernel> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 template <> const char *get_spec_constant_symbolic_ID<SpecConst1>() {
diff --git a/sycl/unittests/assert/assert.cpp b/sycl/unittests/assert/assert.cpp
index 94fcdaca7ce75..74bf38d7595bc 100644
--- a/sycl/unittests/assert/assert.cpp
+++ b/sycl/unittests/assert/assert.cpp
@@ -50,6 +50,7 @@ template <> struct KernelInfo<TestKernel> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 static constexpr const kernel_param_desc_t Signatures[] = {
@@ -68,6 +69,11 @@ struct KernelInfo<::sycl::detail::__sycl_service_kernel__::AssertInfoCopier> {
   static constexpr bool isESIMD() { return 0; }
   static constexpr bool callsThisItem() { return 0; }
   static constexpr bool callsAnyThisFreeFunction() { return 0; }
+  static constexpr int64_t getKernelSize() {
+    // The AssertInfoCopier service kernel lambda captures an accessor.
+    return sizeof(sycl::accessor<sycl::detail::AssertHappened, 1,
+                                 sycl::access::mode::write>);
+  }
 };
 } // namespace detail
 } // namespace sycl
diff --git a/sycl/unittests/buffer/BufferLocation.cpp b/sycl/unittests/buffer/BufferLocation.cpp
index 5d2313ee5dd15..b519447be9e56 100644
--- a/sycl/unittests/buffer/BufferLocation.cpp
+++ b/sycl/unittests/buffer/BufferLocation.cpp
@@ -116,7 +116,8 @@ TEST_F(BufferTest, BufferLocationOnly) {
             cl::sycl::ext::oneapi::accessor_property_list<
                 cl::sycl::ext::intel::property::buffer_location::instance<2>>>
             Acc{Buf, cgh, sycl::read_write, PL};
-        cgh.single_task<TestKernel>([=]() { Acc[0] = 4; });
+        constexpr size_t KS = sizeof(decltype(Acc));
+        cgh.single_task<TestKernel<KS>>([=]() { Acc[0] = 4; });
       })
       .wait();
   EXPECT_EQ(PassedLocation, (uint64_t)2);
@@ -149,7 +150,8 @@ TEST_F(BufferTest, BufferLocationWithAnotherProp) {
                 cl::sycl::ext::intel::property::buffer_location::instance<5>>>
             Acc{Buf, cgh, sycl::write_only, PL};
 
-        cgh.single_task<TestKernel>([=]() { Acc[0] = 4; });
+        constexpr size_t KS = sizeof(decltype(Acc));
+        cgh.single_task<TestKernel<KS>>([=]() { Acc[0] = 4; });
       })
       .wait();
   EXPECT_EQ(PassedLocation, (uint64_t)5);
@@ -209,7 +211,8 @@ TEST_F(BufferTest, WOBufferLocation) {
                        cl::sycl::access::placeholder::false_t,
                        cl::sycl::ext::oneapi::accessor_property_list<>>
             Acc{Buf, cgh, sycl::read_write};
-        cgh.single_task<TestKernel>([=]() { Acc[0] = 4; });
+        constexpr size_t KS = sizeof(decltype(Acc));
+        cgh.single_task<TestKernel<KS>>([=]() { Acc[0] = 4; });
       })
       .wait();
   EXPECT_EQ(PassedLocation, DEFAULT_VALUE);
diff --git a/sycl/unittests/event/EventDestruction.cpp b/sycl/unittests/event/EventDestruction.cpp
index 673c203c4c2be..6b31208cbbba7 100644
--- a/sycl/unittests/event/EventDestruction.cpp
+++ b/sycl/unittests/event/EventDestruction.cpp
@@ -68,11 +68,11 @@ TEST_F(EventDestructionTest, EventDestruction) {
 
     {
       sycl::event E0 = Queue.submit([&](cl::sycl::handler &cgh) {
-        cgh.single_task<TestKernel>([]() {});
+        cgh.single_task<TestKernel<>>([]() {});
       });
       E1 = Queue.submit([&](cl::sycl::handler &cgh) {
         cgh.depends_on(E0);
-        cgh.single_task<TestKernel>([]() {});
+        cgh.single_task<TestKernel<>>([]() {});
       });
       E1.wait();
     }
@@ -85,7 +85,7 @@ TEST_F(EventDestructionTest, EventDestruction) {
 
     sycl::event E2 = Queue.submit([&](cl::sycl::handler &cgh) {
       cgh.depends_on(E1);
-      cgh.single_task<TestKernel>([]() {});
+      cgh.single_task<TestKernel<>>([]() {});
     });
     E2.wait();
     // Dependencies of E1 should be cleared here. It depends on E0.
@@ -93,7 +93,7 @@ TEST_F(EventDestructionTest, EventDestruction) {
 
     sycl::event E3 = Queue.submit([&](cl::sycl::handler &cgh) {
       cgh.depends_on({E1, E2});
-      cgh.single_task<TestKernel>([]() {});
+      cgh.single_task<TestKernel<>>([]() {});
     });
     E3.wait();
     // Dependency of E1 has already cleared. E2 depends on E1 that
@@ -107,20 +107,20 @@ TEST_F(EventDestructionTest, EventDestruction) {
     sycl::buffer<int, 1> Buf(&data[0], sycl::range<1>(2));
     Queue.submit([&](cl::sycl::handler &cgh) {
       auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
-      cgh.single_task<TestKernel>([=]() {});
+      cgh.single_task<TestKernel<>>([=]() {});
     });
 
     Queue.submit([&](cl::sycl::handler &cgh) {
       auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
-      cgh.single_task<TestKernel>([=]() {});
+      cgh.single_task<TestKernel<>>([=]() {});
     });
     sycl::event E1 = Queue.submit([&](cl::sycl::handler &cgh) {
       auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
-      cgh.single_task<TestKernel>([=]() {});
+      cgh.single_task<TestKernel<>>([=]() {});
     });
     sycl::event E2 = Queue.submit([&](cl::sycl::handler &cgh) {
       auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
-      cgh.single_task<TestKernel>([=]() {});
+      cgh.single_task<TestKernel<>>([=]() {});
     });
     E2.wait();
     // Dependencies are deleted through one level of dependencies. When
@@ -172,11 +172,11 @@ TEST_F(EventDestructionTest, GetWaitList) {
 
     {
       sycl::event E0 = Queue.submit([&](cl::sycl::handler &cgh) {
-        cgh.single_task<TestKernel>([]() {});
+        cgh.single_task<TestKernel<>>([]() {});
       });
       E1 = Queue.submit([&](cl::sycl::handler &cgh) {
         cgh.depends_on(E0);
-        cgh.single_task<TestKernel>([]() {});
+        cgh.single_task<TestKernel<>>([]() {});
       });
       E1.wait();
       auto wait_list = E1.get_wait_list();
@@ -190,13 +190,13 @@ TEST_F(EventDestructionTest, GetWaitList) {
 
     sycl::event E2 = Queue.submit([&](cl::sycl::handler &cgh) {
       cgh.depends_on(E1);
-      cgh.single_task<TestKernel>([]() {});
+      cgh.single_task<TestKernel<>>([]() {});
     });
     E2.wait();
 
     sycl::event E3 = Queue.submit([&](cl::sycl::handler &cgh) {
       cgh.depends_on({E1, E2});
-      cgh.single_task<TestKernel>([]() {});
+      cgh.single_task<TestKernel<>>([]() {});
     });
     E3.wait();
 
diff --git a/sycl/unittests/helpers/TestKernel.hpp b/sycl/unittests/helpers/TestKernel.hpp
index d8c2a8b87caec..fd3bd7511668e 100644
--- a/sycl/unittests/helpers/TestKernel.hpp
+++ b/sycl/unittests/helpers/TestKernel.hpp
@@ -10,12 +10,12 @@
 
 #include "PiImage.hpp"
 
-class TestKernel;
+template <size_t KernelSize = 1> class TestKernel;
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
-template <> struct KernelInfo<TestKernel> {
+template <size_t KernelSize> struct KernelInfo<TestKernel<KernelSize>> {
   static constexpr unsigned getNumParams() { return 0; }
   static const kernel_param_desc_t &getParamDesc(int) {
     static kernel_param_desc_t Dummy;
@@ -25,6 +25,7 @@ template <> struct KernelInfo<TestKernel> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return KernelSize; }
 };
 
 } // namespace detail
diff --git a/sycl/unittests/kernel-and-program/Cache.cpp b/sycl/unittests/kernel-and-program/Cache.cpp
index 5a41da137ed28..694576070bdc5 100644
--- a/sycl/unittests/kernel-and-program/Cache.cpp
+++ b/sycl/unittests/kernel-and-program/Cache.cpp
@@ -46,6 +46,7 @@ struct MockKernelInfo {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 template <> struct KernelInfo<TestKernel> : public MockKernelInfo {
diff --git a/sycl/unittests/kernel-and-program/MultipleDevsCache.cpp b/sycl/unittests/kernel-and-program/MultipleDevsCache.cpp
index 15ac42424d728..5a56a999b8e2e 100644
--- a/sycl/unittests/kernel-and-program/MultipleDevsCache.cpp
+++ b/sycl/unittests/kernel-and-program/MultipleDevsCache.cpp
@@ -151,11 +151,12 @@ TEST_F(MultipleDeviceCacheTest, ProgramRetain) {
 
     auto Bundle = cl::sycl::get_kernel_bundle<sycl::bundle_state::input>(
         Queue.get_context());
-    Queue.submit(
-        [&](cl::sycl::handler &cgh) { cgh.single_task<TestKernel>([]() {}); });
+    Queue.submit([&](cl::sycl::handler &cgh) {
+      cgh.single_task<TestKernel<>>([]() {});
+    });
 
     auto BundleObject = cl::sycl::build(Bundle, Bundle.get_devices());
-    auto KernelID = cl::sycl::get_kernel_id<TestKernel>();
+    auto KernelID = cl::sycl::get_kernel_id<TestKernel<>>();
     auto Kernel = BundleObject.get_kernel(KernelID);
 
     // Because of emulating 2 devices program is retained for each one in
diff --git a/sycl/unittests/misc/KernelBuildOptions.cpp b/sycl/unittests/misc/KernelBuildOptions.cpp
index 6ca5967420336..7ad5f993820c1 100644
--- a/sycl/unittests/misc/KernelBuildOptions.cpp
+++ b/sycl/unittests/misc/KernelBuildOptions.cpp
@@ -33,6 +33,7 @@ template <> struct KernelInfo<TestKernel> {
   static constexpr bool isESIMD() { return true; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 } // namespace detail
diff --git a/sycl/unittests/program_manager/EliminatedArgMask.cpp b/sycl/unittests/program_manager/EliminatedArgMask.cpp
index 1e83d0d52abdd..8c9d7c75fce96 100644
--- a/sycl/unittests/program_manager/EliminatedArgMask.cpp
+++ b/sycl/unittests/program_manager/EliminatedArgMask.cpp
@@ -20,8 +20,8 @@
 
 class EAMTestKernel;
 class EAMTestKernel2;
-const char EAMTestKernelName[] = "EAMTestKernel";
-const char EAMTestKernel2Name[] = "EAMTestKernel2";
+constexpr const char EAMTestKernelName[] = "EAMTestKernel";
+constexpr const char EAMTestKernel2Name[] = "EAMTestKernel2";
 constexpr unsigned EAMTestKernelNumArgs = 4;
 
 __SYCL_INLINE_NAMESPACE(cl) {
@@ -37,6 +37,7 @@ template <> struct KernelInfo<EAMTestKernel> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 template <> struct KernelInfo<EAMTestKernel2> {
@@ -49,6 +50,7 @@ template <> struct KernelInfo<EAMTestKernel2> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 } // namespace detail
diff --git a/sycl/unittests/program_manager/SubDevices.cpp b/sycl/unittests/program_manager/SubDevices.cpp
index 72227506f4aee..b0ff4f6a78f01 100644
--- a/sycl/unittests/program_manager/SubDevices.cpp
+++ b/sycl/unittests/program_manager/SubDevices.cpp
@@ -145,11 +145,11 @@ TEST(SubDevices, DISABLED_BuildProgramForSubdevices) {
   sycl::detail::ProgramManager::getInstance().getBuiltPIProgram(
       sycl::detail::OSUtil::getOSModuleHandle(&devBin),
       sycl::detail::getSyclObjImpl(Ctx), subDev1,
-      sycl::detail::KernelInfo<TestKernel>::getName());
+      sycl::detail::KernelInfo<TestKernel<>>::getName());
   // This call should re-use built binary from the cache. If piProgramBuild is
   // called again, the test will fail as second call of redefinedProgramBuild
   sycl::detail::ProgramManager::getInstance().getBuiltPIProgram(
       sycl::detail::OSUtil::getOSModuleHandle(&devBin),
       sycl::detail::getSyclObjImpl(Ctx), subDev2,
-      sycl::detail::KernelInfo<TestKernel>::getName());
+      sycl::detail::KernelInfo<TestKernel<>>::getName());
 }
diff --git a/sycl/unittests/program_manager/itt_annotations.cpp b/sycl/unittests/program_manager/itt_annotations.cpp
index ad024bf43501d..eefa9478ab5ee 100644
--- a/sycl/unittests/program_manager/itt_annotations.cpp
+++ b/sycl/unittests/program_manager/itt_annotations.cpp
@@ -234,7 +234,7 @@ TEST(ITTNotify, UseKernelBundle) {
   auto ExecBundle = sycl::build(KernelBundle);
   Queue.submit([&](sycl::handler &CGH) {
     CGH.use_kernel_bundle(ExecBundle);
-    CGH.single_task<TestKernel>([] {}); // Actual kernel does not matter
+    CGH.single_task<TestKernel<>>([] {}); // Actual kernel does not matter
   });
 
   EXPECT_EQ(HasITTEnabled, true);
@@ -275,7 +275,7 @@ TEST(ITTNotify, VarNotSet) {
   auto ExecBundle = sycl::build(KernelBundle);
   Queue.submit([&](sycl::handler &CGH) {
     CGH.use_kernel_bundle(ExecBundle);
-    CGH.single_task<TestKernel>([] {}); // Actual kernel does not matter
+    CGH.single_task<TestKernel<>>([] {}); // Actual kernel does not matter
   });
 
   EXPECT_EQ(HasITTEnabled, false);
diff --git a/sycl/unittests/program_manager/passing_link_and_compile_options.cpp b/sycl/unittests/program_manager/passing_link_and_compile_options.cpp
index b50f3ae5cf807..d732d474f5ee6 100644
--- a/sycl/unittests/program_manager/passing_link_and_compile_options.cpp
+++ b/sycl/unittests/program_manager/passing_link_and_compile_options.cpp
@@ -37,6 +37,7 @@ template <> struct KernelInfo<EAMTestKernel1> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 template <> struct KernelInfo<EAMTestKernel2> {
@@ -49,6 +50,7 @@ template <> struct KernelInfo<EAMTestKernel2> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 } // namespace detail
diff --git a/sycl/unittests/queue/GetProfilingInfo.cpp b/sycl/unittests/queue/GetProfilingInfo.cpp
index a51b6ef9c30b0..2cefa047edde1 100644
--- a/sycl/unittests/queue/GetProfilingInfo.cpp
+++ b/sycl/unittests/queue/GetProfilingInfo.cpp
@@ -34,6 +34,7 @@ template <> struct KernelInfo<InfoTestKernel> {
   static constexpr bool isESIMD() { return false; }
   static constexpr bool callsThisItem() { return false; }
   static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
 };
 
 } // namespace detail
diff --git a/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp b/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp
index 99a12427c5e7a..dbe604944da43 100644
--- a/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp
+++ b/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp
@@ -59,7 +59,7 @@ TEST_F(SchedulerTest, InOrderQueueHostTaskDeps) {
 
   event Evt = InOrderQueue.submit([&](sycl::handler &CGH) {
     CGH.use_kernel_bundle(ExecBundle);
-    CGH.single_task<TestKernel>([] {});
+    CGH.single_task<TestKernel<>>([] {});
   });
   InOrderQueue
       .submit([&](sycl::handler &CGH) {
diff --git a/sycl/unittests/scheduler/RequiredWGSize.cpp b/sycl/unittests/scheduler/RequiredWGSize.cpp
index 9c5a667626cfb..39c6964f9e653 100644
--- a/sycl/unittests/scheduler/RequiredWGSize.cpp
+++ b/sycl/unittests/scheduler/RequiredWGSize.cpp
@@ -221,7 +221,7 @@ static void performChecks() {
   auto ExecBundle = sycl::build(KernelBundle);
   Queue.submit([&](sycl::handler &CGH) {
     CGH.use_kernel_bundle(ExecBundle);
-    CGH.single_task<TestKernel>([] {}); // Actual kernel does not matter
+    CGH.single_task<TestKernel<>>([] {}); // Actual kernel does not matter
   });
 
   EXPECT_EQ(KernelGetGroupInfoCalled, true);
diff --git a/sycl/unittests/stream/stream.cpp b/sycl/unittests/stream/stream.cpp
index e230939ec7057..ce755bfd08aa9 100644
--- a/sycl/unittests/stream/stream.cpp
+++ b/sycl/unittests/stream/stream.cpp
@@ -73,7 +73,7 @@ TEST(Stream, TestStreamConstructorExceptionNoAllocation) {
       FAIL() << "Unexpected exception was thrown.";
     }
 
-    CGH.single_task<TestKernel>([=]() {});
+    CGH.single_task<TestKernel<>>([=]() {});
   });
 
   ASSERT_EQ(GBufferCreateCounter, 0u) << "Buffers were unexpectedly created.";