diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 63eb048212776..5c42709930436 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/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( + MQueue->getHandleRef(), RawEvents.size(), + RawEvents.size() ? &RawEvents[0] : nullptr, &Event); + MEvent->setHandle(Event); + return Result; + } } return UR_RESULT_ERROR_INVALID_OPERATION; } diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index a7ac73f9e4c34..1daa436b1ee87 100644 --- a/sycl/source/handler.cpp +++ b/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(); - MLastEvent = detail::createSyclObjFromImpl(Event); - return MLastEvent; - } + CommandGroup.reset(new detail::CG(detail::CGType::None, + std::move(impl->CGData), MCodeLoc)); break; } diff --git a/sycl/test-e2e/Basic/empty_command.cpp b/sycl/test-e2e/Basic/empty_command.cpp index 313ca81743c36..dac5865ae8d72 100644 --- a/sycl/test-e2e/Basic/empty_command.cpp +++ b/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); }