[SYCL] Honor dependencies of empty command groups (intel#16180)

Co-authored-by: Sergey Semenov <[email protected]>

Author: aelovikov-intel

sycl/source/detail/scheduler/commands.cpp

@@ -3592,10 +3592,21 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
 
     return UR_RESULT_SUCCESS;
   }
-  case CGType::None:
-    throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
-                          "CG type not implemented. " +
-                              codeToString(UR_RESULT_ERROR_INVALID_OPERATION));
+  case CGType::None: {
+    if (RawEvents.empty()) {
+      // urEnqueueEventsWait with zero events acts like a barrier which is NOT
+      // what we want here. On the other hand, there is nothing to wait for, so
+      // we don't need to enqueue anything.
+      return UR_RESULT_SUCCESS;
+    }
+    const detail::AdapterPtr &Adapter = MQueue->getAdapter();
+    ur_event_handle_t Event;
+    ur_result_t Result = Adapter->call_nocheck<UrApiKind::urEnqueueEventsWait>(
+        MQueue->getHandleRef(), RawEvents.size(),
+        RawEvents.size() ? &RawEvents[0] : nullptr, &Event);
+    MEvent->setHandle(Event);
+    return Result;
+  }
   }
   return UR_RESULT_ERROR_INVALID_OPERATION;
 }

sycl/source/handler.cpp

@@ -496,21 +496,8 @@ event handler::finalize() {
         MCodeLoc));
     break;
   case detail::CGType::None:
-    if (detail::ur::trace(detail::ur::TraceLevel::TRACE_ALL)) {
-      std::cout << "WARNING: An empty command group is submitted." << std::endl;
-    }
-
-    // Empty nodes are handled by Graph like standard nodes
-    // For Standard mode (non-graph),
-    // empty nodes are not sent to the scheduler to save time
-    if (impl->MGraph || (MQueue && MQueue->getCommandGraph())) {
-      CommandGroup.reset(new detail::CG(detail::CGType::None,
-                                        std::move(impl->CGData), MCodeLoc));
-    } else {
-      detail::EventImplPtr Event = std::make_shared<sycl::detail::event_impl>();
-      MLastEvent = detail::createSyclObjFromImpl<event>(Event);
-      return MLastEvent;
-    }
+    CommandGroup.reset(new detail::CG(detail::CGType::None,
+                                      std::move(impl->CGData), MCodeLoc));
     break;
   }
 

sycl/test-e2e/Basic/empty_command.cpp

@@ -26,18 +26,11 @@ void test_host_task_dep() {
   auto empty_cg_event =
       q.submit([&](handler &cgh) { cgh.depends_on(host_event); });
 
-  // FIXME: This should deadlock, but the dependency is ignored currently.
-  empty_cg_event.wait();
-
   assert(x == 0);
   start_execution.count_down();
 
   empty_cg_event.wait();
-  // FIXME: uncomment once the bug mentioned above is fixed.
-  // assert(x == 42);
-
-  // I'm seeing some weird hang without this:
-  host_event.wait();
+  assert(x == 42);
 }
 
 void test_device_event_dep() {
@@ -53,17 +46,12 @@ void test_device_event_dep() {
   auto empty_cg_event =
       q.submit([&](handler &cgh) { cgh.depends_on(device_event); });
 
-  // FIXME: This should deadlock, but the dependency is ignored currently.
-  empty_cg_event.wait();
-
   assert(*p == 0);
   start_execution.count_down();
 
   empty_cg_event.wait();
-  // FIXME: uncomment once the bug mentioned above is fixed.
-  // assert(*p == 42);
+  assert(*p == 42);
 
-  q.wait();
   sycl::free(p, q);
 }
 
@@ -90,17 +78,12 @@ void test_accessor_dep() {
   auto empty_cg_event =
       q.submit([&](handler &cgh) { sycl::accessor a{b, cgh}; });
 
-  // FIXME: This should deadlock, but the dependency is ignored currently.
-  empty_cg_event.wait();
-
   assert(*p == 0);
   start_execution.count_down();
 
   empty_cg_event.wait();
-  // FIXME: uncomment once the bug mentioned above is fixed.
-  // assert(*p == 42);
+  assert(*p == 42);
 
-  q.wait();
   sycl::free(p, q);
 }