cudaDeviceSynchronize 生成新线程，即使设置为阻塞

Question

我正在尝试编写一些代码来确保所有 GPU 活动（尤其是所有正在运行的线程）都停止。我需要这样做以使用 dlclose 卸载模块，因此我需要确保主机和设备上的所有线程都已停止。

根据CUDA documentation，cudaDeviceSynchronize：

阻塞直到设备完成所有先前请求的任务...如果为该设备设置了 cudaDeviceScheduleBlockingSync 标志，主机线程将阻塞直到设备完成其工作。

但是，当我设置阻塞同步标志并调用 cudaDeviceSynchronize 时，会产生一个新的主机线程，它在 cudaDeviceSynchronize 返回后仍在运行。这与我想要实现的目标相反。

此行为在示例程序中进行了演示：

#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;
}

如果我用nvcc -g main.cu 编译这个程序，并在gdb 中运行它，调用info threads 表明在cudaDeviceSynchronize 返回后有两个线程在运行。

在gdb中运行时cudaDeviceSynchronise后一行信息线程的输出：

(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

谁能帮我理解为什么 cudaDeviceSynchronize 会产生一个新线程，以及为什么该线程在调用返回后仍在运行？

谁能指出我正确的方向，帮助我找到一种阻止方法，直到所有设备和主机活动/线程都完成？

Answer 1

CUDA 4.2 及更高版本具有中间工作线程，用于调解应用程序线程和操作系统之间的阻塞调用。我的测试表明，为您的应用程序使用的每个 GPU 创建一个线程（每个 CUDA 上下文一个线程？）。我怀疑引入这些工作线程是为了使流事件回调的实现更容易（我认为这些线程可能会执行回调）；不过，在这个技术原因上，我可能完全错了。

我真的希望 NVIDIA 提供一个环境变量来禁用这些中间线程。如果您想将程序作为 SCHED_FIFO 运行，则会出现问题。在调用任何 CUDA 例程之前，您必须确保转换到 SCHED_FIFO。否则，在 SCHED_FIFO 转换之前产生的任何工作线程都将被安排为常规线程，而您的主线程是 SCHED_FIFO。这会导致优先级反转，您的主线程被阻塞，等待以较低优先级调度工作线程。在任何线程产生之前转换到 SCHED_FIFO 允许未来的线程继承父级的 SCHED_FIFO 策略和优先级。

至于您的问题的解决方案：您可以在应用程序的上下文中调用 cudaDeviceReset() 吗？这应该有望重新初始化系统中的任何 CUDA 运行时状态并终止任何工作线程。否则，总会有 pthread_cancel()（或 Windows 等效函数），但这可能会使 CUDA 处于不一致的状态。