cudaWaitExternalSemaphoresAsync blocks CPU Kernel launch

[matan4442]

Hi, I am using Vulkan timeline semaphores in CUDA for CPU-GPU synchronization.
My code is:

        for (size_t i = 0; i < num_layers; i++)
        {
            waitVulkanSemaphore(i, ioStream1);
            cudaEventRecord(events_get[i], ioStream1); 
        }

        for (size_t i = 0; i < num_layers; i++)
        {
            cudaStreamWaitEvent(computeStream, events_get[i], 0);
            mockKernel<<<numBlocks, blockSize, 0, computeStream>>>(d_data, dataSize, 42);
        }


        for (size_t i = 0; i <= num_layers; i++)
        {
            // sleep mocks CPU work - can be in different thread
            sleep_ms(600);
            signalVulkanSemaphore(i);
        }

The problem is that if I call waitVulkanSemaphore (calls to cudaWaitExternalSemaphoresAsync ), the call to mockKernel blocks the CPU. Then i don’t get to signalVulkanSemaphore, and i get a deadlock. Why the call to cudaWaitExternalSemaphoresAsync makes the mockKernel to be synchronous and not asynchronous? The full code of is attached.

Thanks very much.
vulkanCode.txt (11.9 KB)

[nv-lduc]

Hi @matan4442,

It seems like the code you posted here has a bug. External synchronization objects do not support the “wait-before-signal” behavior, as explained in the programming guide here: 1. Introduction — CUDA C++ Programming Guide

See this section in particular:

Synchronization objects can be imported into CUDA using cudaImportExternalSemaphore(). An imported synchronization object can then be signaled using cudaSignalExternalSemaphoresAsync() and waited on using cudaWaitExternalSemaphoresAsync(). It is illegal to issue a wait before the corresponding signal has been issued.

Vulkan timeline semaphores do support wait-before-signal, but that capability does not translate to Cuda when importing one.

If you are looking for a solution, you may have more luck discussing this particular use-case in the CUDA developer forums.

I hope that helps,

Lionel

[matan4442]

Thx @nv-lduc
I see..
why would I want to use semaphores without wait before signal anyway? if I know that a semaphore must be signaled before I wait on it, doesn’t it misses all the point of waiting on it at all?