diff --git a/sycl/source/detail/graph_impl.hpp b/sycl/source/detail/graph_impl.hpp
index 21c44159fceb..fce836ea19e0 100644
--- a/sycl/source/detail/graph_impl.hpp
+++ b/sycl/source/detail/graph_impl.hpp
@@ -374,6 +374,16 @@ class graph_impl {
     if (PropList.has_property<property::graph::no_cycle_check>()) {
       MSkipCycleChecks = true;
     }
+    if (SyclDevice.get_info<
+            ext::oneapi::experimental::info::device::graph_support>() ==
+        info::graph_support_level::unsupported) {
+      std::stringstream Stream;
+      Stream << SyclDevice.get_backend();
+      std::string BackendString = Stream.str();
+      throw sycl::exception(
+          sycl::make_error_code(errc::invalid),
+          BackendString + " backend is not supported by SYCL Graph extension.");
+    }
   }
 
   /// Remove node from list of root nodes.
diff --git a/sycl/test-e2e/Graph/exception_unsupported_backend.cpp b/sycl/test-e2e/Graph/exception_unsupported_backend.cpp
new file mode 100644
index 000000000000..8b3e079aa476
--- /dev/null
+++ b/sycl/test-e2e/Graph/exception_unsupported_backend.cpp
@@ -0,0 +1,35 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Tests the ability to finalize a empty command graph
+// The test checks that invalid exception is thrown
+// when trying to create a graph with an unsupported backend.
+
+#include "graph_common.hpp"
+
+int GetUnsupportedBackend(const sycl::device &Dev) {
+  // Return 1 if the device backend is unsupported or 0 else.
+  // 0 does not prevent another device to be picked as a second choice
+  return Dev.get_info<
+             ext::oneapi::experimental::info::device::graph_support>() ==
+         ext::oneapi::experimental::info::graph_support_level::unsupported;
+}
+
+int main() {
+  sycl::device Dev{GetUnsupportedBackend};
+  queue Queue{Dev};
+
+  if (Dev.get_info<ext::oneapi::experimental::info::device::graph_support>() !=
+      ext::oneapi::experimental::info::graph_support_level::unsupported)
+    return 0;
+
+  std::error_code ExceptionCode = make_error_code(sycl::errc::success);
+  try {
+    exp_ext::command_graph Graph{Queue.get_context(), Dev};
+  } catch (exception &Exception) {
+    ExceptionCode = Exception.code();
+  }
+  assert(ExceptionCode == sycl::errc::invalid);
+
+  return 0;
+}
diff --git a/sycl/unittests/Extensions/CommandGraph.cpp b/sycl/unittests/Extensions/CommandGraph.cpp
index beacadce8bc0..3d0a8ed5bb9a 100644
--- a/sycl/unittests/Extensions/CommandGraph.cpp
+++ b/sycl/unittests/Extensions/CommandGraph.cpp
@@ -15,6 +15,7 @@
 #include <detail/config.hpp>
 #include <helpers/PiMock.hpp>
 #include <helpers/ScopedEnvVar.hpp>
+#include <helpers/TestKernel.hpp>
 
 #include <gtest/gtest.h>
 
@@ -394,7 +395,7 @@ TEST_F(CommandGraphTest, AddNode) {
   ASSERT_TRUE(GraphImpl->MRoots.empty());
 
   auto Node1 = Graph.add(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
   ASSERT_NE(sycl::detail::getSyclObjImpl(Node1), nullptr);
   ASSERT_FALSE(sycl::detail::getSyclObjImpl(Node1)->isEmpty());
   ASSERT_EQ(GraphImpl->MRoots.size(), 1lu);
@@ -449,17 +450,17 @@ TEST_F(CommandGraphTest, Finalize) {
   sycl::buffer<int> Buf(1);
   auto Node1 = Graph.add([&](sycl::handler &cgh) {
     sycl::accessor A(Buf, cgh, sycl::write_only, sycl::no_init);
-    cgh.single_task<class TestKernel1>([=]() { A[0] = 1; });
+    cgh.single_task<TestKernel<>>([]() {});
   });
 
   // Add independent node
   auto Node2 = Graph.add(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
 
   // Add a node that depends on Node1 due to the accessor
   auto Node3 = Graph.add([&](sycl::handler &cgh) {
     sycl::accessor A(Buf, cgh, sycl::write_only, sycl::no_init);
-    cgh.single_task<class TestKernel2>([=]() { A[0] = 3; });
+    cgh.single_task<TestKernel<>>([]() {});
   });
 
   // Guarantee order of independent nodes 1 and 2
@@ -485,7 +486,7 @@ TEST_F(CommandGraphTest, MakeEdge) {
 
   // Add two independent nodes
   auto Node1 = Graph.add(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
   auto Node2 = Graph.add([&](sycl::handler &cgh) {});
   ASSERT_EQ(GraphImpl->MRoots.size(), 2ul);
   ASSERT_TRUE(sycl::detail::getSyclObjImpl(Node1)->MSuccessors.empty());
@@ -579,7 +580,7 @@ TEST_F(CommandGraphTest, BeginEndRecording) {
 TEST_F(CommandGraphTest, GetCGCopy) {
   auto Node1 = Graph.add([&](sycl::handler &cgh) {});
   auto Node2 = Graph.add(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); },
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); },
       {experimental::property::node::depends_on(Node1)});
 
   // Get copy of CG of Node2 and check equality
@@ -601,21 +602,21 @@ TEST_F(CommandGraphTest, GetCGCopy) {
 TEST_F(CommandGraphTest, SubGraph) {
   // Add sub-graph with two nodes
   auto Node1Graph = Graph.add(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
   auto Node2Graph = Graph.add(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); },
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); },
       {experimental::property::node::depends_on(Node1Graph)});
   auto GraphExec = Graph.finalize();
 
   // Add node to main graph followed by sub-graph and another node
   experimental::command_graph MainGraph(Queue.get_context(), Dev);
   auto Node1MainGraph = MainGraph.add(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
   auto Node2MainGraph =
       MainGraph.add([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); },
                     {experimental::property::node::depends_on(Node1MainGraph)});
   auto Node3MainGraph = MainGraph.add(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); },
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); },
       {experimental::property::node::depends_on(Node2MainGraph)});
 
   // Assert order of the added sub-graph
@@ -653,10 +654,10 @@ TEST_F(CommandGraphTest, RecordSubGraph) {
   // Record sub-graph with two nodes
   Graph.begin_recording(Queue);
   auto Node1Graph = Queue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
   auto Node2Graph = Queue.submit([&](sycl::handler &cgh) {
     cgh.depends_on(Node1Graph);
-    cgh.single_task<class TestKernel>([]() {});
+    cgh.single_task<TestKernel<>>([]() {});
   });
   Graph.end_recording(Queue);
   auto GraphExec = Graph.finalize();
@@ -665,14 +666,14 @@ TEST_F(CommandGraphTest, RecordSubGraph) {
   experimental::command_graph MainGraph(Queue.get_context(), Dev);
   MainGraph.begin_recording(Queue);
   auto Node1MainGraph = Queue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
   auto Node2MainGraph = Queue.submit([&](handler &cgh) {
     cgh.depends_on(Node1MainGraph);
     cgh.ext_oneapi_graph(GraphExec);
   });
   auto Node3MainGraph = Queue.submit([&](sycl::handler &cgh) {
     cgh.depends_on(Node2MainGraph);
-    cgh.single_task<class TestKernel>([]() {});
+    cgh.single_task<TestKernel<>>([]() {});
   });
   MainGraph.end_recording(Queue);
 
@@ -722,7 +723,7 @@ TEST_F(CommandGraphTest, InOrderQueue) {
   // Record in-order queue with three nodes
   InOrderGraph.begin_recording(InOrderQueue);
   auto Node1Graph = InOrderQueue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
 
   auto PtrNode1 =
       sycl::detail::getSyclObjImpl(InOrderGraph)
@@ -731,7 +732,7 @@ TEST_F(CommandGraphTest, InOrderQueue) {
   ASSERT_TRUE(PtrNode1->MPredecessors.empty());
 
   auto Node2Graph = InOrderQueue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
 
   auto PtrNode2 =
       sycl::detail::getSyclObjImpl(InOrderGraph)
@@ -744,7 +745,7 @@ TEST_F(CommandGraphTest, InOrderQueue) {
   ASSERT_EQ(PtrNode2->MPredecessors.front().lock(), PtrNode1);
 
   auto Node3Graph = InOrderQueue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
 
   auto PtrNode3 =
       sycl::detail::getSyclObjImpl(InOrderGraph)
@@ -782,7 +783,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmpty) {
   // node
   InOrderGraph.begin_recording(InOrderQueue);
   auto Node1Graph = InOrderQueue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
 
   auto PtrNode1 =
       sycl::detail::getSyclObjImpl(InOrderGraph)
@@ -803,7 +804,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmpty) {
   ASSERT_EQ(PtrNode2->MPredecessors.front().lock(), PtrNode1);
 
   auto Node3Graph = InOrderQueue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
 
   auto PtrNode3 =
       sycl::detail::getSyclObjImpl(InOrderGraph)
@@ -847,7 +848,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyFirst) {
   ASSERT_TRUE(PtrNode1->MPredecessors.empty());
 
   auto Node2Graph = InOrderQueue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
 
   auto PtrNode2 =
       sycl::detail::getSyclObjImpl(InOrderGraph)
@@ -860,7 +861,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyFirst) {
   ASSERT_EQ(PtrNode2->MPredecessors.front().lock(), PtrNode1);
 
   auto Node3Graph = InOrderQueue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
 
   auto PtrNode3 =
       sycl::detail::getSyclObjImpl(InOrderGraph)
@@ -896,7 +897,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyLast) {
   // Record in-order queue with two regular nodes then an empty node
   InOrderGraph.begin_recording(InOrderQueue);
   auto Node1Graph = InOrderQueue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
 
   auto PtrNode1 =
       sycl::detail::getSyclObjImpl(InOrderGraph)
@@ -905,7 +906,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyLast) {
   ASSERT_TRUE(PtrNode1->MPredecessors.empty());
 
   auto Node2Graph = InOrderQueue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
 
   auto PtrNode2 =
       sycl::detail::getSyclObjImpl(InOrderGraph)
@@ -947,9 +948,9 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyLast) {
 TEST_F(CommandGraphTest, MakeEdgeErrors) {
   // Set up some nodes in the graph
   auto NodeA = Graph.add(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
   auto NodeB = Graph.add(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
 
   // Test error on calling make_edge when a queue is recording to the graph
   Graph.begin_recording(Queue);
@@ -982,7 +983,7 @@ TEST_F(CommandGraphTest, MakeEdgeErrors) {
   experimental::command_graph<experimental::graph_state::modifiable> GraphOther{
       Queue.get_context(), Queue.get_device()};
   auto NodeOther = GraphOther.add(
-      [&](sycl::handler &cgh) { cgh.single_task<class TestKernel>([]() {}); });
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
 
   ASSERT_THROW(
       {
diff --git a/sycl/unittests/helpers/PiMockPlugin.hpp b/sycl/unittests/helpers/PiMockPlugin.hpp
index 23d1f8f24daa..dac3fce0fddc 100644
--- a/sycl/unittests/helpers/PiMockPlugin.hpp
+++ b/sycl/unittests/helpers/PiMockPlugin.hpp
@@ -163,7 +163,7 @@ inline pi_result mock_piDeviceGetInfo(pi_device device,
                                       size_t *param_value_size_ret) {
   constexpr char MockDeviceName[] = "Mock device";
   constexpr char MockSupportedExtensions[] =
-      "cl_khr_fp64 cl_khr_fp16 cl_khr_il_program";
+      "cl_khr_fp64 cl_khr_fp16 cl_khr_il_program ur_exp_command_buffer";
   switch (param_name) {
   case PI_DEVICE_INFO_TYPE: {
     // Act like any device is a GPU.