diff --git a/sycl/include/sycl/detail/kernel_desc.hpp b/sycl/include/sycl/detail/kernel_desc.hpp
index 92bb17b9b4c7a..e3134accc29f2 100644
--- a/sycl/include/sycl/detail/kernel_desc.hpp
+++ b/sycl/include/sycl/detail/kernel_desc.hpp
@@ -188,94 +188,6 @@ template <typename KNT> struct KernelIdentity {
   using type = KNT;
 };
 
-// KernelRegistrar registers a kernel for later lookup and production of a
-// kernel_id via sycl::get_kernel_id<KN>(). Registration is accomplished by
-// instantiating a specialization with the kernel name type and KernelInfo
-// for each kernel and then calling the registerKernelName() static
-// member function to provoke an ODR-use of the getKernelNameHelper()
-// hidden friend function definition. This makes the definition of
-// getKernelNameHelper() for a particular kernel name available for use
-// by arbitrary translation units without requiring registration in each
-// of them.
-template <typename KN, typename KernelInfo> struct KernelRegistrar {
-  // getKernelNameHelper() provides an interface to retrieve a kernel
-  // entry point name corresponding to a kernel name type. Though dependent
-  // on the KN template parameter, this function is not a member of the
-  // KernelRegistrar class; it is a namespace scope function that can be
-  // redeclared in other class template contexts. Since it is only dependent
-  // on KN, note that multiple instantiations of KernelRegistrar with the
-  // same kernel name but different kernel entry point name symbol would
-  // result in duplicate definitions of getKernelNameHelper(). This is
-  // by design.
-  //
-  // Hidden friend function definitions are implicitly inline. C++17
-  // [basic.def.odr]p4 states:
-  //   "... An inline function or variable shall be defined in every
-  //   translation unit in which it is odr-used outside of a discarded
-  //   statement."
-  // This requirement is intentionally violated in order to provoke a
-  // linker error when a call to sycl::get_kernel_id<KN>() is made for
-  // a type that has not been registered as a kernel name type; see
-  // the comments associated with the KernelRegistry class definition.
-  // The quoted statement allows implementations to not emit definitions
-  // of inline functions even if ODR-used within a translation unit (e.g.,
-  // because all such uses were inlined). In practice, some implementations
-  // do skip emitting such definitions subject to compile-time and link-time
-  // optimization. The GNU 'used' attribute is applied to force GCC and
-  // Clang to emit such definitions. MSVC does not appear to suppress such
-  // definitions, even with whole program analysis (/WP) and link time
-  // code generation (/LTCG).
-  //
-  // Name lookup for getKernelNameHelper() requires argument dependent
-  // lookup (ADL); even from within this class definition. The call to
-  // getKernelNameHelper() in registerKernelName() therefore passes an
-  // instance of this class (dependent on both KN and KernelInfo) to enable
-  // lookup in this class. A conversion operator then enables the argument
-  // to be converted to the KernelIdentity wrapper type (only dependent on KN).
-#if defined(__GNUG__) || defined(__clang__)
-  __attribute__((used))
-#endif
-#if defined(_WIN32)
-  __declspec(dllexport)
-#endif
-  friend const char *
-  getKernelNameHelper(KernelIdentity<KN>) noexcept {
-    return KernelInfo::getName();
-  }
-  operator KernelIdentity<KN>() const noexcept { return KernelIdentity<KN>{}; }
-  static void registerKernelName() noexcept {
-    (void)getKernelNameHelper(KernelRegistrar<KN, KernelInfo>{});
-  }
-};
-
-// KernelRegistry enables retrieval of a kernel entry point name given
-// kernel name type KN. Lookup is performed by instantiating a
-// specialization with the kernel name type and then calling the
-// getKernelName() static member function. getKernelName() then calls
-// getKernelNameHelper(); a friend function that is only defined if
-// type KN was registered as a kernel name type in some translation unit
-// using the KernelRegistrar class template above. As with KernelRegistrar,
-// the call to getKernelNameHelper() requires ADL, so an instance of this
-// class is passed as an argument which is convertible to the kernel
-// name wrapper type.
-template <typename KN> struct KernelRegistry {
-#if defined(__GNUC__) && !defined(__clang__)
-#pragma GCC diagnostic push
-#pragma GCC diagnostic ignored "-Wnon-template-friend"
-#endif
-#if defined(_WIN32)
-  __declspec(dllexport)
-#endif
-  friend const char *getKernelNameHelper(KernelIdentity<KN>) noexcept;
-#if defined(__GNUC__) && !defined(__clang__)
-#pragma GCC diagnostic pop
-#endif
-  operator KernelIdentity<KN>() const noexcept { return KernelIdentity<KN>{}; }
-  static const char *getKernelName() noexcept {
-    return getKernelNameHelper(KernelRegistry<KN>{});
-  }
-};
-
 } // namespace detail
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 38ba8f872e5af..ed0fa77820e37 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -799,8 +799,7 @@ class __SYCL_EXPORT handler {
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
-    detail::KernelRegistrar<NameT,
-                            detail::KernelInfo<NameT>>::registerKernelName();
+
 #ifndef __SYCL_DEVICE_ONLY__
     throwIfActionIsCreated();
     throwOnKernelParameterMisuse(Info);
@@ -962,8 +961,6 @@ class __SYCL_EXPORT handler {
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     (void)Props;
     constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
-    detail::KernelRegistrar<NameT,
-                            detail::KernelInfo<NameT>>::registerKernelName();
     detail::KernelWrapper<WrapAsVal, NameT, KernelType, ElementType,
                           PropertiesT>::wrap(KernelFunc);
 
diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp
index b15736d304c94..933449aeab625 100644
--- a/sycl/include/sycl/kernel_bundle.hpp
+++ b/sycl/include/sycl/kernel_bundle.hpp
@@ -602,8 +602,10 @@ __SYCL_EXPORT kernel_id get_kernel_id_impl(string_view KernelName);
 
 /// \returns the kernel_id associated with the KernelName
 template <typename KernelName> kernel_id get_kernel_id() {
+  // FIXME: This must fail at link-time if KernelName not in any available
+  // translation units.
   return detail::get_kernel_id_impl(
-      std::string_view(detail::KernelRegistry<KernelName>::getKernelName()));
+      detail::CompileTimeKernelInfo<KernelName>.Name);
 }
 
 /// \returns a vector with all kernel_id's defined in the application
diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp
index df41891758768..629b73e8258f3 100644
--- a/sycl/include/sycl/queue.hpp
+++ b/sycl/include/sycl/queue.hpp
@@ -3916,8 +3916,7 @@ auto submit_kernel_direct(const queue &Queue,
   detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
 
   using KernelType = std::decay_t<KernelTypeUniversalRef>;
-  detail::KernelRegistrar<KernelName,
-                          detail::KernelInfo<KernelName>>::registerKernelName();
+
   detail::KernelWrapper<WrapAs, KernelName, KernelType, LambdaArgType,
                         PropertiesT>::wrap(KernelFunc);
 
diff --git a/sycl/source/detail/device_kernel_info.cpp b/sycl/source/detail/device_kernel_info.cpp
index dcc4d3ed9649d..c86f411ae1a42 100644
--- a/sycl/source/detail/device_kernel_info.cpp
+++ b/sycl/source/detail/device_kernel_info.cpp
@@ -50,17 +50,7 @@ void DeviceKernelInfo::setCompileTimeInfoIfNeeded(
   if (!isCompileTimeInfoSet())
     CompileTimeKernelInfoTy::operator=(Info);
   assert(isCompileTimeInfoSet());
-// FIXME On Windows, if we have a duplicate kernel name in multiple DSOs, the
-// full equality assertion can be violated. In this case, we would still be
-// reusing the kernel associated with this device kernel info (on Linux or
-// Windows), and while we could issue a diagnostic for the Windows case here,
-// we don't in order to maintain parity with Linux behavior, which happens
-// to work if the kernels are identical.
-#ifdef _WIN32
-  assert(std::string_view{Info.Name} == std::string_view{this->Name});
-#else
   assert(Info == *this);
-#endif
 }
 
 void DeviceKernelInfo::setImplicitLocalArgPos(int Pos) {
diff --git a/sycl/test-e2e/KernelAndProgram/kernel-bundle-across-tus.cpp b/sycl/test-e2e/KernelAndProgram/kernel-bundle-across-tus.cpp
deleted file mode 100644
index f8dc349f0b667..0000000000000
--- a/sycl/test-e2e/KernelAndProgram/kernel-bundle-across-tus.cpp
+++ /dev/null
@@ -1,59 +0,0 @@
-// RUN: %{build} -DMAIN -c -o %t.main.o
-// RUN: %{build} -c -o %t.kernel.o
-// RUN: %clangxx -fsycl %{sycl_target_opts} %t.main.o %t.kernel.o -o %t.out
-// RUN: %{run} %t.out
-
-// DEFINE: %{dynamic_lib_suffix} = %if windows %{dll%} %else %{so%}
-// RUN: rm -rf %t.dir; mkdir -p %t.dir
-// RUN: %clangxx -fsycl %{sycl_target_opts} %fPIC %shared_lib %s -o %t.dir/libkernel.%{dynamic_lib_suffix}
-// RUN: %if !windows %{%{run-aux}%} \
-// RUN: %clangxx -fsycl %{sycl_target_opts} %t.main.o -o %t.dir/%{t:stem}.out -L%t.dir \
-// RUN: %if windows                                                                    \
-// RUN:   %{%t.dir/libkernel.lib%}                                                     \
-// RUN: %else                                                                          \
-// RUN:   %{-L%t.dir -lkernel -Wl,-rpath=%t.dir%}
-
-// RUN: %{run} %t.dir/%{t:stem}.out
-#include <sycl/detail/core.hpp>
-#include <sycl/kernel_bundle.hpp>
-
-#if defined(_WIN32)
-#define API_EXPORT __declspec(dllexport)
-#else
-#define API_EXPORT
-#endif
-
-class KernelA;
-API_EXPORT void submitKernel(sycl::queue &Q);
-#ifndef MAIN
-void submitKernel(sycl::queue &Q) {
-  Q.submit([&](sycl::handler &CGH) { CGH.single_task<KernelA>([=]() {}); });
-}
-#else
-int main() {
-  sycl::queue Queue;
-  sycl::device Dev = Queue.get_device();
-  sycl::context Ctx = Queue.get_context();
-  sycl::kernel_bundle KernelBundle =
-      sycl::get_kernel_bundle<sycl::bundle_state::executable>(Ctx, {Dev});
-
-  auto FoundKernel = KernelBundle.get_kernel<KernelA>();
-  auto ExpectedKernel = KernelBundle.get_kernel(sycl::get_kernel_id<KernelA>());
-
-  // operator== isn't guaranteed to work in this scenario, so compare traits
-  // about the kernels instead.
-  auto FoundKernelName =
-      FoundKernel.get_info<sycl::info::kernel::function_name>();
-  auto ExpectedKernelName =
-      ExpectedKernel.get_info<sycl::info::kernel::function_name>();
-
-  auto FoundKernelContext = FoundKernel.get_info<sycl::info::kernel::context>();
-  auto ExpectedKernelContext =
-      ExpectedKernel.get_info<sycl::info::kernel::context>();
-
-  assert(FoundKernelName == ExpectedKernelName);
-  assert(FoundKernelContext == ExpectedKernelContext);
-
-  return 0;
-}
-#endif
diff --git a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/a.hpp b/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/a.hpp
deleted file mode 100644
index 36cfb49e2eb9e..0000000000000
--- a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/a.hpp
+++ /dev/null
@@ -1,7 +0,0 @@
-#include <sycl/detail/core.hpp>
-
-#include "common.hpp"
-
-class TestKernel;
-API_EXPORT void submitKernelDirectA(sycl::queue &Q, int *Ptr);
-API_EXPORT void submitKernelWithIdA(sycl::queue &Q, int *Ptr);
diff --git a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/a_kernel.cpp b/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/a_kernel.cpp
deleted file mode 100644
index 994185d1ff187..0000000000000
--- a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/a_kernel.cpp
+++ /dev/null
@@ -1,7 +0,0 @@
-#include "a.hpp"
-
-void submitKernelDirectA(sycl::queue &Q, int *Ptr) {
-  Q.submit([&](sycl::handler &CGH) {
-    CGH.single_task<TestKernel>([=]() { Ptr[0] = 1; });
-  });
-}
diff --git a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/a_kernel_id.cpp b/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/a_kernel_id.cpp
deleted file mode 100644
index 3a100477cb4a0..0000000000000
--- a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/a_kernel_id.cpp
+++ /dev/null
@@ -1,5 +0,0 @@
-#include "a.hpp"
-
-void submitKernelWithIdA(sycl::queue &Q, int *Ptr) {
-  enqueueWithKernelId(Q, sycl::get_kernel_id<TestKernel>(), Ptr);
-}
diff --git a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/b.hpp b/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/b.hpp
deleted file mode 100644
index db3c315aaadbf..0000000000000
--- a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/b.hpp
+++ /dev/null
@@ -1,7 +0,0 @@
-#include <sycl/detail/core.hpp>
-
-#include "common.hpp"
-
-class TestKernel;
-API_EXPORT void submitKernelDirectB(sycl::queue &Q, int *Ptr);
-API_EXPORT void submitKernelWithIdB(sycl::queue &Q, int *Ptr);
diff --git a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/b_kernel.cpp b/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/b_kernel.cpp
deleted file mode 100644
index cdc797ab30111..0000000000000
--- a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/b_kernel.cpp
+++ /dev/null
@@ -1,7 +0,0 @@
-#include "b.hpp"
-
-void submitKernelDirectB(sycl::queue &Q, int *Ptr) {
-  Q.submit([&](sycl::handler &CGH) {
-    CGH.single_task<TestKernel>([=]() { Ptr[0] = 2; });
-  });
-}
diff --git a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/b_kernel_id.cpp b/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/b_kernel_id.cpp
deleted file mode 100644
index ad1a96e245377..0000000000000
--- a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/b_kernel_id.cpp
+++ /dev/null
@@ -1,5 +0,0 @@
-#include "b.hpp"
-
-void submitKernelWithIdB(sycl::queue &Q, int *Ptr) {
-  enqueueWithKernelId(Q, sycl::get_kernel_id<TestKernel>(), Ptr);
-}
diff --git a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/common.hpp b/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/common.hpp
deleted file mode 100644
index 5c5bf29403218..0000000000000
--- a/sycl/test-e2e/SharedLib/DuplicateKernelName/Inputs/common.hpp
+++ /dev/null
@@ -1,21 +0,0 @@
-#pragma once
-
-#include <sycl/detail/core.hpp>
-#include <sycl/kernel_bundle.hpp>
-
-#if defined(_WIN32)
-#define API_EXPORT __declspec(dllexport)
-#else
-#define API_EXPORT
-#endif
-
-inline void enqueueWithKernelId(sycl::queue &Q, sycl::kernel_id Id, int *Ptr) {
-  sycl::kernel_bundle KernelBundle =
-      sycl::get_kernel_bundle<sycl::bundle_state::executable>(Q.get_context(),
-                                                              {Q.get_device()});
-  auto Kernel = KernelBundle.get_kernel(Id);
-  Q.submit([&](sycl::handler &CGH) {
-    CGH.set_args(Ptr);
-    CGH.single_task(Kernel);
-  });
-}
diff --git a/sycl/test-e2e/SharedLib/DuplicateKernelName/duplicate_kernel_name.cpp b/sycl/test-e2e/SharedLib/DuplicateKernelName/duplicate_kernel_name.cpp
deleted file mode 100644
index ab3541aa09c7b..0000000000000
--- a/sycl/test-e2e/SharedLib/DuplicateKernelName/duplicate_kernel_name.cpp
+++ /dev/null
@@ -1,55 +0,0 @@
-
-// REQUIRES: aspect-usm_shared_allocations
-
-// RUN: rm -rf %t.dir; mkdir -p %t.dir
-
-// DEFINE: %{dynamic_lib_suffix} = %if windows %{dll%} %else %{so%}
-// RUN: %clangxx -fsycl %{sycl_target_opts} %fPIC %shared_lib %S/Inputs/a_kernel_id.cpp %S/Inputs/a_kernel.cpp -o %t.dir/liba.%{dynamic_lib_suffix}
-// RUN: %clangxx -fsycl %{sycl_target_opts} %fPIC %shared_lib %S/Inputs/b_kernel_id.cpp %S/Inputs/b_kernel.cpp -o %t.dir/libb.%{dynamic_lib_suffix}
-// RUN: %if !windows %{%{run-aux}%} \
-// RUN: %clangxx -fsycl %{sycl_target_opts} %s -o %t.dir/%{t:stem}.out -L%t.dir \
-// RUN: %if windows                                                                    \
-// RUN:   %{%t.dir/liba.lib %t.dir/libb.lib%}                                                     \
-// RUN: %else                                                                          \
-// RUN:   %{-L%t.dir -la -lb -Wl,-rpath=%t.dir%}
-
-// RUN: %{run} %t.dir/%{t:stem}.out
-#include <sycl/usm.hpp>
-
-#include "Inputs/a.hpp"
-#include "Inputs/b.hpp"
-#include "Inputs/common.hpp"
-
-using namespace sycl;
-
-// This test simply captures the current way we handle duplicate kernel names
-// in multiiple DSOs: there's no diagnostic issued for this, and the backend
-// will pick one of the two images. There's no guarantee the backend will pick
-// the same image every time we need to select one, so while direct submissions
-// should reuse the same kernel via caching, there's no guarantee which kernel
-// ends up being submitted using the kernel bundle API.
-
-int main() {
-  queue Q;
-  int *Ptr = malloc_shared<int>(1, Q);
-
-  auto RunTest = [&](auto Func) {
-    *Ptr = 0;
-    Func(Q, Ptr);
-    Q.wait();
-    assert(*Ptr == 1 || *Ptr == 2);
-  };
-
-  RunTest(submitKernelDirectA);
-  int FirstResult = *Ptr;
-  RunTest(submitKernelDirectB);
-  assert(*Ptr == FirstResult);
-  RunTest([](queue &Q, int *Ptr) {
-    enqueueWithKernelId(Q, sycl::get_kernel_id<TestKernel>(), Ptr);
-  });
-  RunTest(submitKernelWithIdA);
-  RunTest(submitKernelWithIdB);
-
-  free(Ptr, Q);
-  return 0;
-}
diff --git a/sycl/test/basic_tests/get_invalid_kernel_id.cpp b/sycl/test/basic_tests/get_invalid_kernel_id.cpp
deleted file mode 100644
index e7e47a4af20e5..0000000000000
--- a/sycl/test/basic_tests/get_invalid_kernel_id.cpp
+++ /dev/null
@@ -1,11 +0,0 @@
-// RUN: not %clangxx -fsycl %s 2>&1 | FileCheck %s --check-prefix %if windows %{CHECK-WINDOWS%} %else %{CHECK-LINUX%}
-
-#include <sycl/sycl.hpp>
-
-class NotAKernel;
-
-int main() {
-  // CHECK-LINUX: {{undefined reference to|undefined symbol:}} {{.?}}sycl::_V1::detail::getKernelNameHelper(sycl::_V1::detail::KernelIdentity<NotAKernel>)
-  // CHECK-WINDOWS: unresolved external symbol "char const * __cdecl sycl::_V1::detail::getKernelNameHelper(struct sycl::_V1::detail::KernelIdentity<class NotAKernel>)"
-  sycl::get_kernel_id<NotAKernel>();
-}
diff --git a/sycl/unittests/Extensions/RootGroup.cpp b/sycl/unittests/Extensions/RootGroup.cpp
index 517d385c9a33b..6ddbb67a8d0ad 100644
--- a/sycl/unittests/Extensions/RootGroup.cpp
+++ b/sycl/unittests/Extensions/RootGroup.cpp
@@ -43,8 +43,6 @@ struct KernelInfo<QueryKernel> : public unittest::MockKernelInfoBase {
 } // namespace _V1
 } // namespace sycl
 
-template void sycl::unittest::registerKernelNames<QueryKernel>();
-
 static sycl::unittest::MockDeviceImage Img =
     sycl::unittest::generateDefaultImage({"QueryKernel"});
 static const sycl::unittest::MockDeviceImageArray<1> ImgArray{&Img};
diff --git a/sycl/unittests/SYCL2020/KernelID.cpp b/sycl/unittests/SYCL2020/KernelID.cpp
index 28db6cae5c116..ccc1c47292a42 100644
--- a/sycl/unittests/SYCL2020/KernelID.cpp
+++ b/sycl/unittests/SYCL2020/KernelID.cpp
@@ -45,8 +45,6 @@ static sycl::unittest::MockDeviceImage Imgs[2] = {
     sycl::unittest::generateDefaultImage({"KernelID_TestKernel2"})};
 static sycl::unittest::MockDeviceImageArray<2> ImgArray{Imgs};
 
-template void
-sycl::unittest::registerKernelNames<TestKernel1, TestKernel2, TestKernel3>();
 TEST(KernelID, AllProgramKernelIds) {
   std::vector<sycl::kernel_id> AllKernelIDs = sycl::get_kernel_ids();
 
@@ -209,3 +207,18 @@ TEST(KernelID, HasKernelTemplated) {
   EXPECT_FALSE(InputBundle1.has_kernel<TestKernel2>());
   EXPECT_TRUE(InputBundle1.has_kernel<TestKernel3>());
 }
+
+TEST(KernelID, GetKernelIDInvalidKernelName) {
+  sycl::unittest::UrMock<> Mock;
+  sycl::platform Plt = sycl::platform();
+
+  try {
+    sycl::get_kernel_id<class NotAKernel>();
+    FAIL() << "Expected an exception";
+  } catch (sycl::exception const &e) {
+    EXPECT_TRUE(e.code() == sycl::errc::runtime);
+    EXPECT_EQ(std::string("No kernel found with the specified name"), e.what());
+  } catch (...) {
+    FAIL() << "Expected sycl::exception";
+  }
+}
diff --git a/sycl/unittests/helpers/MockKernelInfo.hpp b/sycl/unittests/helpers/MockKernelInfo.hpp
index a308bd2b434cb..b65cf679450e4 100644
--- a/sycl/unittests/helpers/MockKernelInfo.hpp
+++ b/sycl/unittests/helpers/MockKernelInfo.hpp
@@ -32,15 +32,6 @@ struct MockKernelInfoBase {
   static constexpr unsigned getColumnNumber() { return 0; }
 };
 
-// Instantiating this function template registers kernel names in
-// KernelRegistry, as would normally be done by kernel wrappers at their
-// submission.
-template <typename... KernelNames> void registerKernelNames() {
-  (sycl::detail::KernelRegistrar<
-       KernelNames,
-       sycl::detail::KernelInfo<KernelNames>>::registerKernelName(),
-   ...);
-}
 } // namespace unittest
 } // namespace _V1
 } // namespace sycl
@@ -48,8 +39,6 @@ template <typename... KernelNames> void registerKernelNames() {
 // In most cases we don't need to redefine any other method besides getName(),
 // so here we only have the simplest helper. If any test needs to redefine more
 // methods, they can do that explicitly.
-// While kernel name registry is not part of the integration header, all kernels
-// that have one need to have the other, so add the instantiation here.
 #define MOCK_INTEGRATION_HEADER(KernelName)                                    \
   namespace sycl {                                                             \
   inline namespace _V1 {                                                       \
@@ -59,8 +48,5 @@ template <typename... KernelNames> void registerKernelNames() {
     static constexpr const char *getName() { return #KernelName; }             \
   };                                                                           \
   } /* namespace detail */                                                     \
-  namespace unittest {                                                         \
-  template void registerKernelNames<KernelName>();                             \
-  } /* namespace unittest */                                                   \
   } /* namespace _V1 */                                                        \
   } /* namespace sycl */
diff --git a/sycl/unittests/helpers/TestKernel.hpp b/sycl/unittests/helpers/TestKernel.hpp
index 900b0c95e1da5..2bf3e7a043018 100644
--- a/sycl/unittests/helpers/TestKernel.hpp
+++ b/sycl/unittests/helpers/TestKernel.hpp
@@ -81,8 +81,3 @@ static sycl::unittest::MockDeviceImage Imgs[] = {
     sycl::unittest::generateDefaultImage({"TestKernelWithStream"}),
     sycl::unittest::generateDefaultImage({"TestKernelWithPtr"})};
 static sycl::unittest::MockDeviceImageArray<4> ImgArray{Imgs};
-
-inline void registerTestKernelNames() {
-  sycl::unittest::registerKernelNames<
-      TestKernel, TestKernelWithAcc, TestKernelWithStream, TestKernelWithPtr>();
-}
diff --git a/sycl/unittests/kernel-and-program/KernelBuildOptions.cpp b/sycl/unittests/kernel-and-program/KernelBuildOptions.cpp
index d928ec54ae42a..952c71d39bb5a 100644
--- a/sycl/unittests/kernel-and-program/KernelBuildOptions.cpp
+++ b/sycl/unittests/kernel-and-program/KernelBuildOptions.cpp
@@ -32,8 +32,6 @@ struct KernelInfo<BuildOptsTestKernel> : public unittest::MockKernelInfoBase {
 } // namespace _V1
 } // namespace sycl
 
-template void sycl::unittest::registerKernelNames<BuildOptsTestKernel>();
-
 static ur_result_t redefinedProgramBuild(void *pParams) {
   auto params = *static_cast<ur_program_build_exp_params_t *>(pParams);
   if (*params.ppOptions)
diff --git a/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp b/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp
index d9f893eac7053..d2e5d092c1150 100644
--- a/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp
+++ b/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp
@@ -72,8 +72,6 @@ struct KernelInfo<EAMTestKernel3> : public unittest::MockKernelInfoBase {
 } // namespace _V1
 } // namespace sycl
 
-template void sycl::unittest::registerKernelNames<EAMTestKernel, EAMTestKernel2,
-                                                  EAMTestKernel3>();
 static sycl::unittest::MockDeviceImage generateEAMTestKernelImage() {
   using namespace sycl::unittest;
 
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 9e5700994aefc..0ac843e753d41 100644
--- a/sycl/unittests/program_manager/passing_link_and_compile_options.cpp
+++ b/sycl/unittests/program_manager/passing_link_and_compile_options.cpp
@@ -55,10 +55,6 @@ struct KernelInfo<EAMTestKernel3> : public unittest::MockKernelInfoBase {
 } // namespace _V1
 } // namespace sycl
 
-template void
-sycl::unittest::registerKernelNames<EAMTestKernel1, EAMTestKernel2,
-                                    EAMTestKernel3>();
-
 template <typename T>
 static sycl::unittest::MockDeviceImage
 generateEAMTestKernelImage(std::string _cmplOptions, std::string _lnkOptions) {