[SYCL] Implement sycl_khr_queue_empty_query

sycl/include/sycl/queue.hpp

@@ -3486,12 +3486,20 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   /// \return the backend associated with this queue.
   backend get_backend() const noexcept;
 
-  /// Allows to check status of the queue (completed vs noncompleted).
+  /// Allows to check status of the queue (completed vs incomplete).
   ///
   /// \return returns true if all enqueued commands in the queue have been
   /// completed, otherwise returns false.
   bool ext_oneapi_empty() const;
 
+  /// Allows to check status of the queue (completed vs incomplete).
+  ///
+  /// \return returns true if all enqueued commands in the queue have been
+  /// completed, otherwise returns false.
+#ifdef __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
+  bool khr_empty() const;
+#endif
+
   ur_native_handle_t getNative(int32_t &NativeHandleDesc) const;
 
   std::optional<event> ext_oneapi_get_last_event() const {

sycl/source/CMakeLists.txt

@@ -50,6 +50,7 @@ function(add_sycl_rt_library LIB_NAME LIB_OBJ_NAME)
   target_compile_definitions(
     ${LIB_OBJ_NAME}
     PRIVATE
+      __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
       __SYCL_INTERNAL_API
       SYCL2020_DISABLE_DEPRECATION_WARNINGS
       $<$<BOOL:${MSVC}>:__SYCL_BUILD_SYCL_DLL>

sycl/source/detail/queue_impl.cpp

@@ -756,7 +756,7 @@ ur_native_handle_t queue_impl::getNative(int32_t &NativeHandleDesc) const {
   return Handle;
 }
 
-bool queue_impl::ext_oneapi_empty() const {
+bool queue_impl::queue_empty() const {
   // If we have in-order queue where events are not discarded then just check
   // the status of the last event.
   if (isInOrder() && !MDiscardEvents) {

sycl/source/detail/queue_impl.hpp

@@ -593,7 +593,7 @@ class queue_impl {
     MStreamsServiceEvents.push_back(Event);
   }
 
-  bool ext_oneapi_empty() const;
+  bool queue_empty() const;
 
   event memcpyToDeviceGlobal(const std::shared_ptr<queue_impl> &Self,
                              void *DeviceGlobalPtr, const void *Src,

sycl/source/queue.cpp

@@ -392,7 +392,9 @@ bool queue::is_in_order() const {
 
 backend queue::get_backend() const noexcept { return getImplBackend(impl); }
 
-bool queue::ext_oneapi_empty() const { return impl->ext_oneapi_empty(); }
+bool queue::ext_oneapi_empty() const { return impl->queue_empty(); }
+
+bool queue::khr_empty() const { return impl->queue_empty(); }
 
 void queue::ext_oneapi_prod() { impl->flush(); }
 

sycl/test-e2e/Basic/in_order_queue_status_khr_empty.cpp

@@ -0,0 +1,87 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Test checks that queue::khr_empty() returns status of the in-order queue.
+
+#define __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
+
+#include <sycl/detail/core.hpp>
+#include <sycl/properties/all_properties.hpp>
+#include <sycl/usm.hpp>
+
+#include <chrono>
+#include <thread>
+
+static void CheckArray(int *x, size_t buffer_size, int expected) {
+  for (size_t i = 0; i < buffer_size; ++i) {
+    assert(x[i] == expected);
+  }
+}
+
+using namespace sycl;
+
+void TestFunc(queue &Q) {
+  static constexpr int Size = 100;
+
+  assert(Q.khr_empty() && "Queue is expected to be empty");
+
+  int *X = malloc_host<int>(Size, Q);
+  int *Y = malloc_host<int>(Size, Q);
+
+  auto FillEv = Q.fill(X, 99, Size);
+  auto SingleTaskEv = Q.submit([&](handler &CGH) {
+    auto SingleTask = [=] {
+      for (int I = 0; I < Size; I++)
+        X[I] += 1;
+    };
+    CGH.single_task(SingleTask);
+  });
+  auto MemCpyEv = Q.copy(X, Y, Size);
+  constexpr int NumIter = 5;
+  for (int I = 0; I < NumIter; I++) {
+    Q.submit([&](handler &CGH) {
+      CGH.parallel_for<class Kernel1>(sycl::range<1>(Size),
+                                      [=](sycl::id<1> WI) { Y[WI] *= 2; });
+    });
+  }
+
+  // Wait a bit to give a chance for tasks to complete.
+  std::this_thread::sleep_for(std::chrono::milliseconds(500));
+
+  // We expect that all submitted tasks are finished if khr_empty is true.
+  if (Q.khr_empty())
+    CheckArray(Y, Size, 3200);
+
+  Q.wait();
+
+  // After synchronization queue must be empty.
+  assert(Q.khr_empty() && "Queue is expected to be empty");
+
+  free(X, Q);
+  free(Y, Q);
+}
+
+int main() {
+  // Test in-order queue.
+  queue Q1{property::queue::in_order()};
+  TestFunc(Q1);
+
+  // Test in-order queue with discard_events property.
+  sycl::property_list Props{
+      property::queue::in_order{},
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  queue Q2{Props};
+
+  bool ExceptionThrown = false;
+  try {
+    TestFunc(Q2);
+  } catch (sycl::exception &E) {
+    ExceptionThrown = true;
+  }
+
+  // Feature is not supported for OpenCL, exception must be thrown.
+  if (Q2.get_device().get_backend() == backend::opencl)
+    return ExceptionThrown ? 0 : -1;
+
+  return 0;
+}

sycl/test-e2e/Basic/out_of_order_queue_status_khr_empty.cpp

@@ -0,0 +1,82 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Test checks that queue::khr_empty() returns status of the out-of-order
+// queue.
+
+#define __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
+
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#include <chrono>
+#include <thread>
+
+static void CheckArray(int *x, size_t buffer_size, int expected) {
+  for (size_t i = 0; i < buffer_size; ++i) {
+    assert(x[i] == expected);
+  }
+}
+
+using namespace sycl;
+
+void TestFunc(queue &Q) {
+  static constexpr int Size = 100;
+
+  assert(Q.khr_empty() && "Queue is expected to be empty");
+
+  int *X = malloc_host<int>(Size, Q);
+  int *Y = malloc_host<int>(Size, Q);
+
+  auto FillEv = Q.fill(X, 99, Size);
+  auto HostEv = Q.submit([&](handler &CGH) {
+    CGH.depends_on(FillEv);
+    auto HostTask = [=] {
+      for (int I = 0; I < Size; I++)
+        X[I] += 1;
+    };
+    CGH.host_task(HostTask);
+  });
+  auto MemCpyEv = Q.copy(X, Y, Size, {HostEv});
+  constexpr int NumIter = 5;
+  for (int I = 0; I < NumIter; I++) {
+    Q.submit([&](handler &CGH) {
+      CGH.depends_on(MemCpyEv);
+      CGH.parallel_for<class Kernel1>(
+          sycl::range<1>(Size / NumIter),
+          [=](sycl::id<1> WI) { Y[WI + I * Size / NumIter] *= 2; });
+    });
+  }
+
+  // Wait a bit to give a chance for tasks to complete.
+  std::this_thread::sleep_for(std::chrono::milliseconds(500));
+
+  // We expect that all submitted tasks are finished if khr_empty is true.
+  if (Q.khr_empty())
+    CheckArray(Y, Size, 200);
+
+  Q.wait();
+
+  // After synchronization queue must be empty.
+  assert(Q.khr_empty() && "Queue is expected to be empty");
+
+  free(X, Q);
+  free(Y, Q);
+}
+
+int main() {
+  queue Q;
+
+  bool ExceptionThrown = false;
+  try {
+    TestFunc(Q);
+  } catch (sycl::exception &E) {
+    ExceptionThrown = true;
+  }
+
+  // Feature is not supported for OpenCL, exception must be thrown.
+  if (Q.get_device().get_backend() == backend::opencl)
+    return ExceptionThrown ? 0 : -1;
+
+  return 0;
+}

sycl/test/abi/sycl_symbols_linux.dump

@@ -3706,6 +3706,7 @@ _ZNK4sycl3_V15queue8get_infoINS0_4info5queue15reference_countEEENS0_6detail18is_
 _ZNK4sycl3_V15queue8get_infoINS0_4info5queue6deviceEEENS0_6detail18is_queue_info_descIT_E11return_typeEv
 _ZNK4sycl3_V15queue8get_infoINS0_4info5queue7contextEEENS0_6detail18is_queue_info_descIT_E11return_typeEv
 _ZNK4sycl3_V15queue9getNativeERi
+_ZNK4sycl3_V15queue9khr_emptyEv
 _ZNK4sycl3_V16ONEAPI15filter_selector13select_deviceEv
 _ZNK4sycl3_V16ONEAPI15filter_selector5resetEv
 _ZNK4sycl3_V16ONEAPI15filter_selectorclERKNS0_6deviceE

sycl/test/abi/sycl_symbols_windows.dump

@@ -4238,6 +4238,7 @@
 ?is_in_order@queue@_V1@sycl@@QEBA_NXZ
 ?is_specialization_constant_set@kernel_bundle_plain@detail@_V1@sycl@@IEBA_NPEBD@Z
 ?join_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBV?$vector@V?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@V?$allocator@V?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@@2@@5@W4bundle_state@23@@Z
+?khr_empty@queue@_V1@sycl@@QEBA_NXZ
 ?khr_get_default_context@platform@_V1@sycl@@QEBA?AVcontext@23@XZ
 ?lgamma_r_impl@detail@_V1@sycl@@YA?AVhalf@half_impl@123@V45123@PEAH@Z
 ?lgamma_r_impl@detail@_V1@sycl@@YAMMPEAH@Z