Appending a dependency event which is to be signalled from the host causes hangs

[pvelesko]

I create an event which (in the full application) will be signalled from a different thread.
I enqueue a kernel which depends on this event, and when using immediate command lists, the enqueue command itself hangs.

Expected behavior: kernel enqueue returns but the kernel doesn't run until we HostSignal the blocking event.
callbacks.zip

Using regular command lists, I don't see the host hang on the kernel enqueue

[pvelesko]

Driver version 23.17.26241.22
tested on
Intel HD 770
Intel(R) Arc(TM) A380 Graphics

[Sarbojit2019]

@pvelesko,

I looked into the sample (callback.zip) attached here. It looks like there is programming error because of which code is getting into dead lock hence the hang. I have corrected the code in main.cpp file and test passed with immediate command list as well.

Immediate command list tries to complete the dispatch immediatly but in the attached sample it has a wait event hence dispatch call waits for event to complete but event signal call has been kept after it hence dispatch never completes which shows as hang. When I moved launch to separate thread test passed. See the code change below

  L1 // Launch kernel on the GPU
  L2 std::cout << "Launching kernel\n";

  L3 // Launching kernel using worker thread
  L4 // In case of immediate command list it is required to be on other thread else it will deadlock  
  L5 std::thread kernel_launch(kernel_func, cmdList, kernel, &dispatch, Event, 1, &HostSignalEvent);

  L6 //ZE_CHECK(zeCommandListAppendLaunchKernel(cmdList, kernel, &dispatch,
  L7 //                                             Event, 0, nullptr)); //1, &HostSignalEvent));
  L8 std::cout << "Host Signal Blocking Event\n";
  L9 ZE_CHECK(zeEventHostSignal(HostSignalEvent));   

  L10 auto begin = std::chrono::steady_clock::now();
  L11 kernel_launch.join();

L5 where I am launching the kernel in other thread where it will be waiting for event to finish. L9 signals the event which helps kernel_launch to complete.

Let me know if you have any further questions.

[pvelesko]

  cmdQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS;

enumerator ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS
Device execution is scheduled and will complete in future; explicit synchronization object must be used to determine completeness

submitting things to an async command queue should not be blocking operations.

[pjaaskel]

https://spec.oneapi.io/level-zero/latest/core/PROG.html#low-latency-immediate-command-lists: "Commands appended into an immediate command list may execute synchronously, by blocking until the command is complete."

[pjaaskel]

But yep, I'd hope the "blocking possibility" doesn't include also waiting for events the command depends on. It'd cause the cumbersome need to have another thread just to construct the queue in case of the kernel command has event dependencies of any kind? Looks like a spec corner case which we can workaround (if needed) with the separate submit thread as suggested by @Sarbojit2019.

[pvelesko]

I assume it's blocking when a sync queue is requested

…

On Mon, Jul 10, 2023 at 13:48 Pekka Jääskeläinen ***@***.***> wrote: But yep, I'd hope the "blocking possibility" doesn't include also waiting for events the command depends on. It'd cause the cumbersome need to have another thread just to construct the queue in case of the kernel command has event dependencies of any kind? Looks like a spec corner case which we can workaround (if needed) with the separate submit thread as suggested by @Sarbojit2019 <https://github.com/Sarbojit2019>. — Reply to this email directly, view it on GitHub <#658 (comment)>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/ACCJBQPMD5XXJRO6FA7CIY3XPPMY3ANCNFSM6AAAAAAZ3LCYW4> . You are receiving this because you were mentioned.Message ID: ***@***.***>

[pjaaskel]

I assume that it might execute the command before returning (to achieve lowest possible latency), but not deadlock on an input event dep (based on the manual link I posted).

[pvelesko]

Looks like a spec corner case

The spec is ambiguous in this regard. It would make sense if

"Commands appended into an immediate command list may execute synchronously, by blocking until the command is complete."

is true when ZE_COMMAND_QUEUE_MODE_SYNCHRONOUS is set:

"Device execution always completes immediately on execute; Host thread is blocked using wait on implicit synchronization object"

@pjaaskel

which we can workaround (if needed) with the separate submit thread as suggested by

We can just use regular command lists for callbacks to get around this but I'm seeing hangs in the full application and can't reproduce in a simple reproducer. Furthermore, I'm seeing deadlocks on non-blocking operations such as zeEventQuery but I wanted to get clarity on this issue first.

[pjaaskel]

With ZE_COMMAND_QUEUE_MODE_SYNCHRONOUS it says always blocks whereas with immediate lists it says may block, which have a major semantics difference. But anyhow, I'd also not expect blocking in the case of event deps here.

[pvelesko]

and conversely ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS should not be blocking.

"enumerator ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS
Device execution is scheduled and will complete in future; explicit synchronization object must be used to determine completeness"

[pjaaskel]

Yep, I'd expect "async usage" to work in this case even though it might execute the commands in the same thread (sometimes).

[jandres742]

hi

I checked the test and is valid. Test is submitting a kernel for immediate submission in asynchronous way. This means upon returning from appendLaunchKernel, the kernel has been submitted, and then asynchronously will wait for the event, but the host thread calling for appendLaunchKernel shouldn't block.

Bacdtrace on GPU MAX 1550 shows we are stuck on the submission path:

#0  0x00007ffff7b3fcab in sched_yield () at ../sysdeps/unix/syscall-template.S:120
#1  0x00007ffff6a5a4cc in __gthread_yield () at /usr/include/x86_64-linux-gnu/c++/10/bits/gthr-default.h:693
#2  std::this_thread::yield () at /usr/include/c++/10/thread:379
#3  NEO::WaitUtils::waitFunctionWithPredicate<unsigned long>(unsigned long const volatile*, unsigned long, std::function<bool (unsigned long, unsigned long)>) (predicate=..., expectedValue=1, pollAddress=0xfffffff7fc5000)
    at ../shared/source/utilities/wait_util.h:33
#4  NEO::WaitUtils::waitFunction (expectedValue=1, pollAddress=0xfffffff7fc5000) at ../shared/source/utilities/wait_util.h:38
#5  NEO::CommandStreamReceiver::baseWaitFunction (this=0x55555624d6d0, pollAddress=0xfffffff7fc5000, params=..., taskCountToWait=1) at ../shared/source/command_stream/command_stream_receiver.cpp:445
#6  0x00007ffff6762def in L0::CommandListCoreFamilyImmediate<(GFXCORE_FAMILY)3080>::executeCommandListImmediateWithFlushTaskImpl (this=this@entry=0x55555627b5f0, performMigration=<optimized out>, performMigration@entry=true,
    hasStallingCmds=hasStallingCmds@entry=true, hasRelaxedOrderingDependencies=hasRelaxedOrderingDependencies@entry=false, cmdQ=<optimized out>) at ../level_zero/core/source/cmdlist/cmdlist_hw_immediate.inl:350
#7  0x00007ffff6762ffa in L0::CommandListCoreFamilyImmediate<(GFXCORE_FAMILY)3080>::executeCommandListImmediateWithFlushTask (this=this@entry=0x55555627b5f0, performMigration=performMigration@entry=true, hasStallingCmds=hasStallingCmds@entry=true,
    hasRelaxedOrderingDependencies=hasRelaxedOrderingDependencies@entry=false) at ../level_zero/core/source/cmdlist/cmdlist_hw_immediate.inl:283
#8  0x00007ffff6763077 in L0::CommandListCoreFamilyImmediate<(GFXCORE_FAMILY)3080>::flushImmediate (this=0x55555627b5f0, inputRet=<optimized out>, performMigration=<optimized out>, hasStallingCmds=<optimized out>,
    hasRelaxedOrderingDependencies=<optimized out>, hSignalEvent=0x555556284410) at ../level_zero/core/source/cmdlist/cmdlist_hw_immediate.inl:844
#9  0x00007ffff66a61ce in L0::zeCommandListAppendLaunchKernel (hCommandList=<optimized out>, kernelHandle=<optimized out>, launchKernelArgs=<optimized out>, hSignalEvent=<optimized out>, numWaitEvents=<optimized out>, phWaitEvents=<optimized out>)
    at ../level_zero/core/source/cmdlist/cmdlist.h:188
#10 0x00007ffff7eea188 in zeCommandListAppendLaunchKernel () from /usr/local/lib/libze_loader.so.1
#11 0x000055555555c3b8 in main (argc=1, argv=0x7fffffffe288) at /home/gta/bin/callbacks/main.cpp:224

The call to waitFunction seems suspicious. If the list is asynchronous, L0 driver shouldn't be waiting.

Test passes with EnableFlushTaskSubmission=0, so it seems we have a bug on that path.

@JablonskiMateusz : Please let me know if an internal tracker has been created for this issue or I can do it.

[pvelesko]

@JablonskiMateusz @jandres742 EnableFlushTaskSubmission=0 does not seem to make any difference. Is this supposed to work on public oneapi releases?

[jandres742]

@pvelesko : that's a debug key, which needs to be set alongside with NEOReadDebugKeys:

export NEOReadDebugKeys=1
export EnableFlushTaskSubmission=0