Bump UR and adjust use of urKernelSuggestMaxCooperativeGroupCountExp (#15966)

Ref: oneapi-src/unified-runtime#2277

---------

Co-authored-by: Martin Morrison-Grant <[email protected]>

Author: igchor

sycl/cmake/modules/UnifiedRuntimeTag.cmake

@@ -1,7 +1,7 @@
-# commit bb64b3e9f6d30b232b204590b94250ce2da6967b
-# Merge: d23739306f2c e625bfdaa99b
+# commit ea0f3a1f5f15f9af7bf40bd13669afeb9ada569c
+# Merge: bb64b3e9f6d3 4a89e1c69a65
 # Author: Martin Grant <[email protected]>
-# Date:   Martin Grant <[email protected]>
-#     Merge pull request #2471 from Bensuo/ben/l0-include-fix
-#     [L0] Fix issue with test include path
-set(UNIFIED_RUNTIME_TAG bb64b3e9f6d30b232b204590b94250ce2da6967b)
+# Date:   Thu Dec 19 11:26:01 2024 +0000
+#     Merge pull request #2277 from igchor/cooperative_fix
+#     [Spec] fix urKernelSuggestMaxCooperativeGroupCountExp
+set(UNIFIED_RUNTIME_TAG ea0f3a1f5f15f9af7bf40bd13669afeb9ada569c)

sycl/source/detail/kernel_impl.hpp

@@ -320,6 +320,7 @@ kernel_impl::queryMaxNumWorkGroups(queue Queue,
   const auto &Adapter = getAdapter();
   const auto &Handle = getHandleRef();
   auto Device = Queue.get_device();
+  auto DeviceHandleRef = sycl::detail::getSyclObjImpl(Device)->getHandleRef();
 
   size_t WG[Dimensions];
   WG[0] = WorkGroupSize[0];
@@ -331,7 +332,8 @@ kernel_impl::queryMaxNumWorkGroups(queue Queue,
   uint32_t GroupCount{0};
   if (auto Result = Adapter->call_nocheck<
                     UrApiKind::urKernelSuggestMaxCooperativeGroupCountExp>(
-          Handle, Dimensions, WG, DynamicLocalMemorySize, &GroupCount);
+          Handle, DeviceHandleRef, Dimensions, WG, DynamicLocalMemorySize,
+          &GroupCount);
       Result != UR_RESULT_ERROR_UNSUPPORTED_FEATURE &&
       Result != UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE) {
     // The feature is supported and the group size is valid. Check for other