diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc
index 86689b94cae69..bd6914f5d4268 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc
@@ -816,7 +816,7 @@ int main() {
     extern "C"
     SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
     void iota(float start, float *ptr) {
-      size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id();
+      size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
       ptr[id] = start + static_cast<float>(id);
     }
   )""";
@@ -845,6 +845,7 @@ int main() {
     sycl::nd_range ndr{{NUM}, {WGSIZE}};
     cgh.parallel_for(ndr, iota);
   }).wait();
+  sycl::free(ptr, q);
 }
 ----
 
@@ -872,13 +873,13 @@ int main() {
 
     SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
     void iota(float start, float *ptr) {
-      size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id();
+      size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
       ptr[id] = start + static_cast<float>(id);
     }
 
-    SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::range_kernel<1>))
+    SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
     void iota(int start, int *ptr) {
-      size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id();
+      size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
       ptr[id] = start + static_cast<int>(id);
     }
   )""";
@@ -910,6 +911,7 @@ int main() {
     sycl::nd_range ndr{{NUM}, {WGSIZE}};
     cgh.parallel_for(ndr, iota);
   }).wait();
+  sycl::free(ptr, q);
 }
 ----
 
diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp
new file mode 100644
index 0000000000000..a2352977f0c2c
--- /dev/null
+++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp
@@ -0,0 +1,78 @@
+//==- kernel_compiler_sycl_jit.cpp --- kernel_compiler extension tests -----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// REQUIRES: (opencl || level_zero)
+// REQUIRES: aspect-usm_device_allocations
+
+// UNSUPPORTED: accelerator
+// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there.
+
+// RUN: %{build} -o %t.out
+// RUN: %{l0_leak_check} %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/usm.hpp>
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+static constexpr size_t NUM = 1024;
+static constexpr size_t WGSIZE = 16;
+
+int main() {
+  sycl::queue q;
+
+  // The source code for a kernel, defined as a SYCL "free function kernel".
+  std::string source = R"""(
+    #include <sycl/sycl.hpp>
+    namespace syclext = sycl::ext::oneapi;
+    namespace syclexp = sycl::ext::oneapi::experimental;
+
+    extern "C"
+    SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+    void iota(float start, float *ptr) {
+      size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+      ptr[id] = start + static_cast<float>(id);
+    }
+  )""";
+
+  // Create a kernel bundle in "source" state.
+  sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
+      syclexp::create_kernel_bundle_from_source(
+          q.get_context(), syclexp::source_language::sycl_jit, source);
+
+  // Compile the kernel.  There is no need to use the "registered_names"
+  // property because the kernel is declared extern "C".
+  sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe =
+      syclexp::build(kb_src);
+
+  // Get the kernel via its compiler-generated name.
+  sycl::kernel iota = kb_exe.ext_oneapi_get_kernel("iota");
+
+  float *ptr = sycl::malloc_shared<float>(NUM, q);
+  q.submit([&](sycl::handler &cgh) {
+     // Set the values of the kernel arguments.
+     cgh.set_args(3.14f, ptr);
+
+     // Launch the kernel according to its type, in this case an nd-range
+     // kernel.
+     sycl::nd_range ndr{{NUM}, {WGSIZE}};
+     cgh.parallel_for(ndr, iota);
+   }).wait();
+
+  constexpr float eps = 0.001;
+  for (int i = 0; i < NUM; i++) {
+    const float truth = 3.14f + static_cast<float>(i);
+    if (std::abs(ptr[i] - truth) > eps) {
+      std::cout << "Result: " << ptr[i] << " expected " << i << "\n";
+      sycl::free(ptr, q);
+      exit(1);
+    }
+  }
+  sycl::free(ptr, q);
+}
diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp
new file mode 100644
index 0000000000000..5bd1fb66f68ed
--- /dev/null
+++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp
@@ -0,0 +1,81 @@
+//==- kernel_compiler_namespaces.cpp --- kernel_compiler extension tests ---==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// REQUIRES: (opencl || level_zero)
+// REQUIRES: aspect-usm_device_allocations
+
+// UNSUPPORTED: accelerator
+// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there.
+
+// RUN: %{build} -o %t.out
+// RUN: %{l0_leak_check} %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/usm.hpp>
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+static constexpr size_t NUM = 1024;
+static constexpr size_t WGSIZE = 16;
+
+int main() {
+  sycl::queue q;
+
+  // The source code for two kernels defined in different namespaces
+  std::string source = R"""(
+    #include <sycl/sycl.hpp>
+    namespace syclext = sycl::ext::oneapi;
+    namespace syclexp = sycl::ext::oneapi::experimental;
+
+    SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+    void iota(int start, int *ptr) {
+      size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+      ptr[id] = start + static_cast<int>(id);
+    }
+
+    namespace mykernels {
+    SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+    void iota(int start, int *ptr) {
+      size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+      ptr[id] = start + static_cast<int>(id);
+    }
+    } // namespace mykernels
+  )""";
+
+  // Create a kernel bundle in "source" state.
+  sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
+      syclexp::create_kernel_bundle_from_source(
+          q.get_context(), syclexp::source_language::sycl_jit, source);
+
+  // Compile the kernel. Select kernel from the mykernels namespace
+  std::string iota_name{"mykernels::iota"};
+  sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe = syclexp::build(
+      kb_src, syclexp::properties{syclexp::registered_names{iota_name}});
+
+  sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name);
+
+  int *ptr = sycl::malloc_shared<int>(NUM, q);
+  q.submit([&](sycl::handler &cgh) {
+     // Set the values of the kernel arguments.
+     cgh.set_args(3, ptr);
+
+     // Launch the kernel according to its type, in this case an nd-range
+     // kernel.
+     sycl::nd_range ndr{{NUM}, {WGSIZE}};
+     cgh.parallel_for(ndr, iota);
+   }).wait();
+
+  for (int i = 0; i < NUM; i++) {
+    if (ptr[i] != i + 3) {
+      std::cout << "Result: " << ptr[i] << " expected " << i << "\n";
+      sycl::free(ptr, q);
+      exit(1);
+    }
+  }
+  sycl::free(ptr, q);
+}
diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp
new file mode 100644
index 0000000000000..51e142f1939db
--- /dev/null
+++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp
@@ -0,0 +1,83 @@
+//==- kernel_compiler_overload.cpp --- kernel_compiler extension tests -----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// REQUIRES: (opencl || level_zero)
+// REQUIRES: aspect-usm_device_allocations
+
+// UNSUPPORTED: accelerator
+// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there.
+
+// RUN: %{build} -o %t.out
+// RUN: %{l0_leak_check} %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/usm.hpp>
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+static constexpr size_t NUM = 1024;
+static constexpr size_t WGSIZE = 16;
+
+int main() {
+  sycl::queue q;
+
+  // The source code for two kernels defined as overloaded functions.
+  std::string source = R"""(
+    #include <sycl/sycl.hpp>
+    namespace syclext = sycl::ext::oneapi;
+    namespace syclexp = sycl::ext::oneapi::experimental;
+
+    SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+    void iota(float start, float *ptr) {
+      size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+      ptr[id] = start + static_cast<float>(id);
+    }
+
+    SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+    void iota(int start, int *ptr) {
+      size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+      ptr[id] = start + static_cast<int>(id);
+    }
+  )""";
+
+  // Create a kernel bundle in "source" state.
+  sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
+      syclexp::create_kernel_bundle_from_source(
+          q.get_context(), syclexp::source_language::sycl_jit, source);
+
+  // Compile the kernel.  Because there are two overloads of "iota", we need to
+  // use a C++ cast to disambiguate between them.  Here, we are selecting the
+  // "int" overload.
+  std::string iota_name{"(void(*)(int, int*))iota"};
+  sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe = syclexp::build(
+      kb_src, syclexp::properties{syclexp::registered_names{iota_name}});
+
+  // Get the kernel by passing the same string we used to construct the
+  // "registered_names" property.
+  sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name);
+
+  int *ptr = sycl::malloc_shared<int>(NUM, q);
+  q.submit([&](sycl::handler &cgh) {
+     // Set the values of the kernel arguments.
+     cgh.set_args(3, ptr);
+
+     // Launch the kernel according to its type, in this case an nd-range
+     // kernel.
+     sycl::nd_range ndr{{NUM}, {WGSIZE}};
+     cgh.parallel_for(ndr, iota);
+   }).wait();
+
+  for (int i = 0; i < NUM; i++) {
+    if (ptr[i] != i + 3) {
+      std::cout << "Result: " << ptr[i] << " expected " << i << "\n";
+      sycl::free(ptr, q);
+      exit(1);
+    }
+  }
+  sycl::free(ptr, q);
+}
diff --git a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp
index dd068fb40752a..23ea51be5d7c4 100644
--- a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp
+++ b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp
@@ -6,7 +6,7 @@
 // CHECK-DAG: README.md
 // CHECK-DAG: lit.cfg.py
 //
-// CHECK-NUM-MATCHES: 7
+// CHECK-NUM-MATCHES: 10
 //
 // This test verifies that `<sycl/sycl.hpp>` isn't used in E2E tests. Instead,
 // fine-grained includes should used, see