在我的运行时cuda程序中，cpu和gpu可以异步计算，但不能协同计算，为什么？答案

【问题标题】：In my cuda program of runtime ,the cpu and gpu can compute Asynchronously,but not cooperatively,why?在我的运行时cuda程序中，cpu和gpu可以异步计算，但不能协同计算，为什么？
【发布时间】：2016-12-05 08:28:08
【问题描述】：

在我的运行时cuda程序中，cpu和gpu可以异步计算，但不能协同计算，为什么？

我测量程序的时间，总时间是cpu计算时间和gpu计算时间的总和。通过视觉轮廓，我发现gpu直到cpu完成才计算。我的目的是 cpu 计算与 gpu 计算同时进行。

平台：窗口 10
cuda 7.5
vs2013

在调试模式下编译的代码（无优化）

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include<time.h>
__global__ void addKernel()
{
    int a ;
    for (int i = 0; i < 10000;i++)
    for (int j = 0; j < 10000;j++)
        a = i;
}
void comput()
{
    int a = 1;
    for (int i = 0; i < 10000;i++)
        for (int j = 0; j < 10000; j++)
        {
            for (int k = 0; k < 100;k++)
                a = j;

        }

}

int main()
{
    cudaSetDevice(0);
    cudaEvent_t start, stop1;
    cudaEventCreate(&start);
    cudaEventCreate(&stop1);

    clock_t ss = clock();
    cudaEventRecord(start,0);   
    addKernel<<<1,64>>>();
    cudaEventRecord(stop1,0);   

    clock_t ct = clock();
    comput();
    clock_t ctt = clock();

    cudaEventSynchronize(stop1);
    cudaDeviceSynchronize();
    clock_t sss = clock();

    float t1;
    cudaEventElapsedTime(&t1, start, stop1);
    printf("clock GPU :%.4f s\n", t1/1000);
    printf("clock cpu:%f s\n",(float) (ctt - ct)/CLOCKS_PER_SEC);
    printf("clock total time: %f s\n", (float)(sss - ss) / CLOCKS_PER_SEC);

    cudaEventDestroy(start);
    cudaEventDestroy(stop1);
    cudaDeviceReset();

}

【问题讨论】：

标签： c++ cuda

【解决方案1】：

这里有几个问题（可能）：

如果您使用的是 WDDM 驱动程序（与 TCC 驱动程序相反），内核启动会分批进行，以减少 WDDM 驱动程序的较高启动开销的影响。这意味着驱动程序将推迟启动addKernel() 等待更多工作，直到遇到cudaEventSynchronize() 调用。¹⁾ 但是此时comput() 已经完成。

因此，在您的示例中，CPU 和 GPU 工作确实不是并行运行的，但 GPU 上的 addKernel() 实际上在 CPU 上运行之后 comput()。
您可以阻止（进一步）批处理并在调用comput() 之前通过inserting a call to cudaStreamQuery(0) 强制立即启动addKernel()。
addKernel() 和 compute() 没有外部可见的效果（它们只设置局部变量 a）并且可能会被编译器完全优化掉。在调试模式下编译可能不会阻止所有这些优化。这会使演示异步执行变得更加困难，因为您只测量内核启动和计时开销。
因此，将它们替换为执行实际工作的代码，例如对向量求和，并将结果存储到全局变量中。
正如 halfelf 在他的回答中指出的那样，分析器可能会在某些条件下同步启动内核。

¹⁾ 如果一段时间内没有进一步的工作，等待也可能超时，addKernel() 可能会在 cudaEventSynchronize() 调用之前启动。

【讨论】：

如你所说，这是驱动问题。我已经通过在调用 comput() 之前插入对 cudaEventQuery(stop1) 的调用解决了这个问题。非常感谢。
是的，cudaEventQuery() 具有立即启动当前批次的相同属性。很高兴这解决了您的问题。

【解决方案2】：

来自cuda programming guide：

程序员可以通过将 CUDA_LAUNCH_BLOCKING 环境变量设置为 1 来为系统上运行的所有 CUDA 应用程序全局禁用内核启动的异步性。此功能仅用于调试目的，不应用作使生产软件可靠运行的方法.

如果硬件计数器是通过分析器（Nsight、Visual Profiler）收集的，则内核启动是同步的，除非启用了并发内核分析。如果异步内存副本涉及未页面锁定的主机内存，它们也将是同步的。

另外如果没有优化，宿主函数会运行很长时间，可能是内核的百万倍。如果经过优化，它实际上不会运行任何东西并立即返回。

我建议尝试分别使用CUDA_LAUNCH_BLOCKING=1 和CUDA_LAUNCH_BLOCKING=0 运行您的二进制文件以测试运行时间。还将您的内核和主机函数修改为有意义的。

【讨论】：