cudaDeviceSynchronise spawns new thread, even when set to blocking

Question

I am trying to write some code to ensure all GPU activity (in particular all running threads) are stopped. I need to do this to unload a module with dlclose, so I need to ensure all threads have stopped on both the host and the device.

According to the CUDA documentation, cudaDeviceSynchronize:

Blocks until the device has completed all preceding requested tasks... If the cudaDeviceScheduleBlockingSync flag was set for this device, the host thread will block until the device has finished its work.

However, when I set the blocking sync flag and call cudaDeviceSynchronize, a new host thread is spawned, which is still running after cudaDeviceSynchronize has returned. This is the opposite of what I am trying to achieve.

This behaviour is demonstrated in an example program:

#include <iostream>

void initialiseDevice()
{
    cudaError result = cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);

    if (cudaSuccess == result)
            std::cout << "Set device flags." << std::endl;
    else
            std::cout << "Could not set device flags. (" << result << ")"
                    << std::endl;
}

void synchroniseDevice()
{
    cudaError result = cudaDeviceSynchronize();

    if (cudaSuccess == result)
            std::cout << "Device synchronise returned success." << std::endl;
    else
            std::cout << "Device synchronise returned error. (" << result << ")"
                    << std::endl;
}

int main()
{
    initialiseDevice();
    sleep(1);
    synchroniseDevice(); // new thread is spawned here
    sleep(1);            // new thread is still running here!
    return 0;
}

If I compile this program with nvcc -g main.cu, and run it in gdb, a call to info threads shows that there are two threads running after cudaDeviceSynchronize has returned.

Output of info threads on the line after cudaDeviceSynchronise when running in gdb:

(gdb) info threads 
  Id   Target Id         Frame 
  2    Thread 0x7ffff5b8b700 (LWP 28458) "a.out" 0x00007ffff75aa023 in select
    () at ../sysdeps/unix/syscall-template.S:82
* 1    Thread 0x7ffff7fd4740 (LWP 28255) "a.out" main () at cuda_test.cu:30

Could anyone help me understand why cudaDeviceSynchronize is spawning a new thread, and why the thread is still running after the call returns?

Could anyone point me in the right direction to help me find a method to block until all device and host activity/threads are finished?

cudaDeviceSynchronize() blocks until all device activity is completed. gdb info threads command is giving information on host activity. cudaDeviceSynchronize() doesn't guarantee anything with respect to host threads. If you want to query device threads the correct command is info cuda threads (in cuda-gdb). One of those threads you have listed appears to be something spun up by a system call. That can happen. The other thread appears to be the one that would have been blocked by cudaDeviceSynchronize(). — Robert Crovella
@RobertCrovella Why didn't you put that on answer instead of comment :)? — dreamcrash
Robert, thanks for your comment. In that case, I suppose I was not quite asking the right question. What I am trying to do is stop all threads on the device, AND the host. cudaDeviceSynchronize may well stop all device threads but it consistently launches a new host thread which continues to run after it has returned. I need all threads to finish. I'll edit the question. — Alex
@alex To stop all host threads, isn't it necessary to terminate the application? — Robert Crovella
@dreamcrash I don't really have a good answer for your question. I'm not always sure what constitutes an answer. In this case, it seems alex's question was actually as he has edited. The question about cudaDeviceSynchronize was ancillary. — Robert Crovella

Glenn Glenn · Accepted Answer · 2013-04-12T21:57:33

CUDA 4.2 and later have intermediary worker threads that mediate blocking calls between application threads and operating system. My testing suggests that one thread gets created for each GPU your application uses (one for each CUDA context?). I suspect these worker threads were introduced to make the implementation of stream event callbacks easier (I think these threads may execute the callbacks); although, I could be entirely wrong on this technical reason.

I really wish NVIDIA would have provided an environment variable to disable these intermediary threads. It introduces problems if you want to run your program as SCHED_FIFO. You must be sure to transition to SCHED_FIFO before any CUDA routines are invoked. Otherwise, any worker threads spawned prior the SCHED_FIFO transition will be scheduled as regular threads while your main thread is SCHED_FIFO. This leads to priority inversions where your main thread is blocked waiting for a worker thread to be scheduled with a lower priority. Transitioning to SCHED_FIFO before any thread spawning allows future threads to inherit the parent's SCHED_FIFO policy and priority.

As for a solution to your problem: Can you call cudaDeviceReset() in the context of your application? This should hopefully reinitialize any CUDA runtime state in your system and kill off any worker threads. Otherwise, there's always pthread_cancel() (or Windows equivalent), but this may leave CUDA in an inconsistent state.

cudaDeviceSynchronise spawns new thread, even when set to blocking

1 Answers