diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt
index a8fbd568bf5fd..2a3ab53851443 100644
--- a/sycl/plugins/unified_runtime/CMakeLists.txt
+++ b/sycl/plugins/unified_runtime/CMakeLists.txt
@@ -95,23 +95,17 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
   endfunction()
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit 396fb20498c315a526c961d7cb645b42795acd2c
-  # Merge: 719bb9cd e2ffea69
+  # commit 905804c2e93dd046140057fd07a5d6191063bedc
+  # Merge: 0a11fb44 d3d3f6e5
   # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-  # Date:   Thu May 23 10:53:03 2024 +0100
-  #     Merge pull request #1501 from RossBrunton/ross/kerneltests
-  #     [Testing] Spec clarifications and testing updates for kernel
-  set(UNIFIED_RUNTIME_TAG 396fb20498c315a526c961d7cb645b42795acd2c)
+  # Date:   Mon May 27 10:34:13 2024 +0100
+  #     Merge pull request #1581 from 0x12CC/l0_cooperative_kernels
+  #     Implement L0 cooperative kernel functions
+  set(UNIFIED_RUNTIME_TAG 905804c2e93dd046140057fd07a5d6191063bedc)
 
   fetch_adapter_source(level_zero
     ${UNIFIED_RUNTIME_REPO}
-    # commit e428745588dd87e5db3c0bc0df1183eb7d0811a5
-    # Merge: fb3cbd16 80d46bb1
-    # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-    # Date:   Fri May 24 10:31:13 2024 +0100
-    #     Merge pull request #1656 from pbalcer/leaks-program-kernel
-    #     fix leaks on level-zero interop program and kernel handles
-    e428745588dd87e5db3c0bc0df1183eb7d0811a5
+    ${UNIFIED_RUNTIME_TAG}
   )
 
   fetch_adapter_source(opencl
diff --git a/sycl/test-e2e/GroupAlgorithm/root_group.cpp b/sycl/test-e2e/GroupAlgorithm/root_group.cpp
index a629a337da802..d7da5455a8a2a 100644
--- a/sycl/test-e2e/GroupAlgorithm/root_group.cpp
+++ b/sycl/test-e2e/GroupAlgorithm/root_group.cpp
@@ -1,5 +1,5 @@
-// Fails with opencl non-cpu and level_zero on linux, enable when fixed.
-// XFAIL: (opencl && !cpu) || (linux && level_zero)
+// Fails with opencl non-cpu, enable when fixed.
+// XFAIL: (opencl && !cpu)
 // RUN: %{build} -I . -o %t.out
 // RUN: %{run} %t.out
 
@@ -48,20 +48,19 @@ void testRootGroup() {
   const auto range = sycl::nd_range<1>{maxWGs * WorkGroupSize, WorkGroupSize};
   q.submit([&](sycl::handler &h) {
     sycl::accessor data{dataBuf, h};
-    h.parallel_for<class RootGroupKernel>(
-        range, props, [=](sycl::nd_item<1> it) {
-          auto root = it.ext_oneapi_get_root_group();
-          data[root.get_local_id()] = root.get_local_id();
-          sycl::group_barrier(root);
+    h.parallel_for<
+        class RootGroupKernel>(range, props, [=](sycl::nd_item<1> it) {
+      auto root = it.ext_oneapi_get_root_group();
+      data[root.get_local_id()] = root.get_local_id();
+      sycl::group_barrier(root);
 
-          root =
-              sycl::ext::oneapi::experimental::this_work_item::get_root_group<
-                  1>();
-          int sum = data[root.get_local_id()] +
-                    data[root.get_local_range() - root.get_local_id() - 1];
-          sycl::group_barrier(root);
-          data[root.get_local_id()] = sum;
-        });
+      root =
+          sycl::ext::oneapi::experimental::this_work_item::get_root_group<1>();
+      int sum = data[root.get_local_id()] +
+                data[root.get_local_range() - root.get_local_id() - 1];
+      sycl::group_barrier(root);
+      data[root.get_local_id()] = sum;
+    });
   });
   sycl::host_accessor data{dataBuf};
   const int workItemCount = static_cast<int>(range.get_global_range().size());
@@ -95,15 +94,15 @@ void testRootGroupFunctions() {
                                  ? root.get_local_id() == sycl::id<1>(0)
                                  : root.get_local_id() == sycl::id<1>(3);
             testResults[2] = root.get_group_range() == sycl::range<1>(1);
-            testResults[3] =
-                root.get_local_range() == sycl::range<1>(WorkGroupSize);
+            testResults[3] = root.get_local_range() == it.get_global_range();
             testResults[4] =
-                root.get_max_local_range() == sycl::range<1>(WorkGroupSize);
+                root.get_max_local_range() == root.get_local_range();
             testResults[5] = root.get_group_linear_id() == 0;
             testResults[6] =
                 root.get_local_linear_id() == root.get_local_id().get(0);
             testResults[7] = root.get_group_linear_range() == 1;
-            testResults[8] = root.get_local_linear_range() == WorkGroupSize;
+            testResults[8] =
+                root.get_local_linear_range() == root.get_local_range().size();
           }
         });
   });