cpu 代码延迟了 cuda 内核执行

Question

我无法理解以下公认的非常简单的代码，它是一个更复杂的项目的简化版本，我现在花了很多时间在其中。 这段代码将在我的系统上运行大约 2000 毫秒。但是当我启用该线路使 cpu 进入睡眠 500 毫秒时，程序将运行更长的时间，大约为 2500 毫秒。

我无法理解这如何符合 cuda 内核相对于主机异步执行的声明？

在 Vistual Studio 2019 上运行 cuda 11.1

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <chrono>
#include <iostream>
#include <numeric>
#include <thread>

__global__ void kernel(double* val, int siz) {
    for (int i = 0; i < siz; i++) val[i] = sqrt(val[i]); //calculate square root for every value in array
}

int main() {
    auto t1 = std::chrono::high_resolution_clock::now();

    const int siz = 1'000'000; //array length
    double* val = new double[siz];
    std::iota(val, val + siz, 0.0); //fill array with 0, 1, 2,...
    double* d_val;

    cudaMalloc(&d_val, sizeof(double) * siz);
    cudaMemcpy(d_val, val, sizeof(double) * siz, cudaMemcpyDefault);
    kernel <<<1, 1 >>> (d_val, siz); //start kernel
    //std::this_thread::sleep_for(std::chrono::milliseconds(500)); //---- putting cpu to sleep also delays kernel execution?
    cudaError_t err = cudaDeviceSynchronize();
    auto t2 = std::chrono::high_resolution_clock::now();

    std::cout << "status: " << cudaGetErrorString(err) << std::endl;
    std::chrono::duration<double, std::milli> ms = t2 - t1;
    std::cout << "duration: " << ms.count() << std::endl;

    delete[] val;
}

Answer 1

我无法理解这如何符合 cuda 内核相对于主机异步执行的说法？

您正在体验 WDDM 命令批处理，如 here 所述。

简而言之，在 Windows 上，当处于 WDDM 驱动程序模型中时，GPU 命令（例如来自 cuda 运行时 API 的任何内容，以及内核启动）将被发送到命令队列。每隔一段时间，根据未发布的启发式，并且没有提供明确的用户控制，命令队列将被“刷新”，即发送到 GPU，此时（如果当前不忙）GPU 将开始处理这些命令。

因此，在 WDDM 设置中，将内核分派到命令队列是非阻塞的（控制立即返回给 CPU 线程）。从命令队列到 GPU 的工作分派遵循其他一些启发式方法。 （内核执行与主机线程异步，无论如何）

如果这是一个问题，您至少有几个选择：

1. 在 Windows 上，切换到 TCC 驱动程序模型中的 GPU。
2. 在 Windows 上，尝试使用链接答案中描述的“黑客”之一。
3. 切换到 linux