[WIP] [Do not merge] Implement sycl_khr_queue_flush

sycl/include/sycl/queue.hpp

@@ -3611,6 +3611,11 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   /// completed, otherwise returns false.
   bool khr_empty() const;
 
+  /// Flushes all commands in the queue to the device, but doesn't wait for them
+  /// to complete unlike wait().
+  ///
+  void khr_flush() const;
+
   std::optional<event> ext_oneapi_get_last_event() const {
     return static_cast<std::optional<event>>(ext_oneapi_get_last_event_impl());
   }

sycl/source/detail/queue_impl.cpp

@@ -1268,6 +1268,15 @@ bool queue_impl::queue_empty() const {
   return IsReady;
 }
 
+void queue_impl::queue_flush() const {
+  if (MGraph.lock()) {
+    throw sycl::exception(make_error_code(errc::invalid),
+                          "flush cannot be called for a queue which is "
+                          "recording to a command graph.");
+  }
+  getAdapter().call<UrApiKind::urKhrFlush>(MQueue);
+}
+
 void queue_impl::revisitUnenqueuedCommandsState(
     const EventImplPtr &CompletedHostTask) {
   if (MIsInorder)

sycl/source/detail/queue_impl.hpp

@@ -613,6 +613,8 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
 
   bool queue_empty() const;
 
+  void queue_flush() const;
+
   EventImplPtr memcpyToDeviceGlobal(void *DeviceGlobalPtr, const void *Src,
                                     bool IsDeviceImageScope, size_t NumBytes,
                                     size_t Offset,

sycl/source/queue.cpp

@@ -247,6 +247,8 @@ bool queue::ext_oneapi_empty() const { return impl->queue_empty(); }
 
 bool queue::khr_empty() const { return impl->queue_empty(); }
 
+void queue::khr_flush() const { return impl->queue_flush(); }
+
 void queue::ext_oneapi_prod() { impl->flush(); }
 
 ur_native_handle_t queue::getNative(int32_t &NativeHandleDesc) const {

unified-runtime/include/unified-runtime/ur_api.h

@@ -512,6 +512,8 @@ typedef enum ur_function_t {
   UR_FUNCTION_QUEUE_GET_GRAPH_EXP = 314,
   /// Enumerator for ::urGraphSetDestructionCallbackExp
   UR_FUNCTION_GRAPH_SET_DESTRUCTION_CALLBACK_EXP = 315,
+  /// Enumerator for ::urKhrFlush
+  UR_FUNCTION_KHR_FLUSH = 316,
   /// @cond
   UR_FUNCTION_FORCE_UINT32 = 0x7fffffff
   /// @endcond
@@ -6607,6 +6609,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueFlush(
     /// [in] handle of the queue to be flushed.
     ur_queue_handle_t hQueue);
 
+UR_APIEXPORT ur_result_t UR_APICALL urKhrFlush(
+    /// [in] handle of the queue to be flushed.
+    ur_queue_handle_t hQueue);
+
 #if !defined(__GNUC__)
 #pragma endregion
 #endif
@@ -14700,6 +14706,14 @@ typedef struct ur_queue_flush_params_t {
   ur_queue_handle_t *phQueue;
 } ur_queue_flush_params_t;
 
+///////////////////////////////////////////////////////////////////////////////
+/// @brief Function parameters for urKhrFlush
+/// @details Each entry is a pointer to the parameter passed to the function;
+///     allowing the callback the ability to modify the parameter's value
+typedef struct ur_khr_flush_params_t {
+  ur_queue_handle_t *phQueue;
+} ur_khr_flush_params_t;
+
 ///////////////////////////////////////////////////////////////////////////////
 /// @brief Function parameters for urQueueBeginGraphCaptureExp
 /// @details Each entry is a pointer to the parameter passed to the function;

unified-runtime/include/unified-runtime/ur_api_funcs.def

@@ -75,6 +75,7 @@ _UR_API(urKernelGetSuggestedLocalWorkSizeWithArgs)
 _UR_API(urKernelSetExecInfo)
 _UR_API(urKernelSetSpecializationConstants)
 _UR_API(urKernelSuggestMaxCooperativeGroupCount)
+_UR_API(urKhrFlush)
 _UR_API(urQueueGetInfo)
 _UR_API(urQueueCreate)
 _UR_API(urQueueRetain)

unified-runtime/include/unified-runtime/ur_ddi.h

@@ -623,6 +623,10 @@ typedef ur_result_t(UR_APICALL *ur_pfnQueueFinish_t)(ur_queue_handle_t);
 /// @brief Function-pointer for urQueueFlush
 typedef ur_result_t(UR_APICALL *ur_pfnQueueFlush_t)(ur_queue_handle_t);
 
+///////////////////////////////////////////////////////////////////////////////
+/// @brief Function-pointer for urQueueFlush
+typedef ur_result_t(UR_APICALL *ur_pfnKhrFlush_t)(ur_queue_handle_t);
+
 ///////////////////////////////////////////////////////////////////////////////
 /// @brief Table of Queue functions pointers
 typedef struct ur_queue_dditable_t {
@@ -634,6 +638,7 @@ typedef struct ur_queue_dditable_t {
   ur_pfnQueueCreateWithNativeHandle_t pfnCreateWithNativeHandle;
   ur_pfnQueueFinish_t pfnFinish;
   ur_pfnQueueFlush_t pfnFlush;
+  ur_pfnKhrFlush_t pfnKhrFlush;
 } ur_queue_dditable_t;
 
 ///////////////////////////////////////////////////////////////////////////////

unified-runtime/source/adapters/cuda/queue.cpp

@@ -171,6 +171,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueFlush(ur_queue_handle_t /*hQueue*/) {
   return UR_RESULT_SUCCESS;
 }
 
+UR_APIEXPORT ur_result_t UR_APICALL urKhrFlush(ur_queue_handle_t /*hQueue*/) {
+  return UR_RESULT_SUCCESS;
+}
+
 UR_APIEXPORT ur_result_t UR_APICALL urQueueGetNativeHandle(
     ur_queue_handle_t hQueue, ur_queue_native_desc_t * /*pDesc*/,
     ur_native_handle_t *phNativeQueue) {

unified-runtime/source/adapters/cuda/ur_interface_loader.cpp

@@ -239,6 +239,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetQueueProcAddrTable(
   pDdiTable->pfnCreateWithNativeHandle = urQueueCreateWithNativeHandle;
   pDdiTable->pfnFinish = urQueueFinish;
   pDdiTable->pfnFlush = urQueueFlush;
+  pDdiTable->pfnKhrFlush = urKhrFlush;
   pDdiTable->pfnGetInfo = urQueueGetInfo;
   pDdiTable->pfnGetNativeHandle = urQueueGetNativeHandle;
   pDdiTable->pfnRelease = urQueueRelease;

unified-runtime/source/adapters/hip/queue.cpp

@@ -196,6 +196,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueFlush(ur_queue_handle_t) {
   return UR_RESULT_SUCCESS;
 }
 
+UR_APIEXPORT ur_result_t UR_APICALL urKhrFlush(ur_queue_handle_t) {
+  return UR_RESULT_SUCCESS;
+}
+
 /// Gets the native HIP handle of a UR queue object
 ///
 /// \param[in] hQueue The UR queue to get the native HIP object of.

unified-runtime/source/adapters/hip/ur_interface_loader.cpp

@@ -239,6 +239,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetQueueProcAddrTable(
   pDdiTable->pfnCreateWithNativeHandle = urQueueCreateWithNativeHandle;
   pDdiTable->pfnFinish = urQueueFinish;
   pDdiTable->pfnFlush = urQueueFlush;
+  pDdiTable->pfnKhrFlush = urKhrFlush;
   pDdiTable->pfnGetInfo = urQueueGetInfo;
   pDdiTable->pfnGetNativeHandle = urQueueGetNativeHandle;
   pDdiTable->pfnRelease = urQueueRelease;

unified-runtime/source/adapters/level_zero/queue.cpp

@@ -929,6 +929,12 @@ ur_result_t urQueueFlush(
   return Queue->executeAllOpenCommandLists();
 }
 
+ur_result_t urKhrFlush(
+    /// [in] handle of the queue to be flushed.
+    ur_queue_handle_t Queue) {
+  return Queue->executeAllOpenCommandLists();
+}
+
 ur_result_t urQueueBeginGraphCaptureExp(ur_queue_handle_t /* hQueue */) {
   UR_LOG_LEGACY(ERR,
                 logger::LegacyMessage("[UR][L0] {} function not implemented!"),

unified-runtime/source/adapters/native_cpu/queue.cpp

@@ -94,6 +94,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueFlush(ur_queue_handle_t /*hQueue*/) {
   DIE_NO_IMPLEMENTATION;
 }
 
+UR_APIEXPORT ur_result_t UR_APICALL urKhrFlush(ur_queue_handle_t /*hQueue*/) {
+
+  DIE_NO_IMPLEMENTATION;
+}
+
 UR_APIEXPORT ur_result_t
 urQueueBeginGraphCaptureExp(ur_queue_handle_t /* hQueue */) {
 

unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp

@@ -239,6 +239,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetQueueProcAddrTable(
   pDdiTable->pfnCreateWithNativeHandle = urQueueCreateWithNativeHandle;
   pDdiTable->pfnFinish = urQueueFinish;
   pDdiTable->pfnFlush = urQueueFlush;
+  pDdiTable->pfnKhrFlush = urKhrFlush;
   pDdiTable->pfnGetInfo = urQueueGetInfo;
   pDdiTable->pfnGetNativeHandle = urQueueGetNativeHandle;
   pDdiTable->pfnRelease = urQueueRelease;

unified-runtime/source/adapters/offload/queue.cpp

@@ -124,6 +124,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueFlush(ur_queue_handle_t) {
   return UR_RESULT_SUCCESS;
 }
 
+UR_APIEXPORT ur_result_t UR_APICALL urKhrFlush(ur_queue_handle_t) {
+  return UR_RESULT_SUCCESS;
+}
+
 UR_APIEXPORT ur_result_t UR_APICALL
 urQueueBeginGraphCaptureExp(ur_queue_handle_t /* hQueue */) {
   return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;

unified-runtime/source/adapters/offload/ur_interface_loader.cpp

@@ -249,6 +249,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetQueueProcAddrTable(
   pDdiTable->pfnCreateWithNativeHandle = urQueueCreateWithNativeHandle;
   pDdiTable->pfnFinish = urQueueFinish;
   pDdiTable->pfnFlush = urQueueFlush;
+  pDdiTable->pfnKhrFlush = urKhrFlush;
   pDdiTable->pfnGetInfo = urQueueGetInfo;
   pDdiTable->pfnGetNativeHandle = urQueueGetNativeHandle;
   pDdiTable->pfnRelease = urQueueRelease;

unified-runtime/source/adapters/opencl/queue.cpp

@@ -300,6 +300,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueFlush(ur_queue_handle_t hQueue) {
   return UR_RESULT_SUCCESS;
 }
 
+UR_APIEXPORT ur_result_t UR_APICALL urKhrFlush(ur_queue_handle_t hQueue) {
+  cl_int RetErr = clFlush(hQueue->CLQueue);
+  CL_RETURN_ON_FAILURE(RetErr);
+  return UR_RESULT_SUCCESS;
+}
+
 UR_APIEXPORT ur_result_t UR_APICALL urQueueRetain(ur_queue_handle_t hQueue) {
   hQueue->RefCount.retain();
   return UR_RESULT_SUCCESS;

unified-runtime/source/adapters/opencl/ur_interface_loader.cpp

@@ -223,6 +223,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetQueueProcAddrTable(
   pDdiTable->pfnCreateWithNativeHandle = urQueueCreateWithNativeHandle;
   pDdiTable->pfnFinish = urQueueFinish;
   pDdiTable->pfnFlush = urQueueFlush;
+  pDdiTable->pfnKhrFlush = urKhrFlush;
   pDdiTable->pfnGetInfo = urQueueGetInfo;
   pDdiTable->pfnGetNativeHandle = urQueueGetNativeHandle;
   pDdiTable->pfnRelease = urQueueRelease;