c++ - 执行 C++ CUDA 内核时黑屏，输出正确结果

Question

我正在使用 CUDA C++ 对数组进行一些简单的并行计算。一切正常，内核输出正确的结果(用串行 CPU 代码检查)，但是当内核正在执行时，我的整个屏幕在内核执行的整个过程中变黑了。我是 CUDA 的新手，所以我可能做错了什么，我似乎无法弄清楚是什么。

#define KERNEL_FOR_ITERS 1e6

__global__ void addKernel(float *c, const float *a, const float *b)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i % 2 == 0)
        for (int j = 0; j < KERNEL_FOR_ITERS; j++)
            c[i] += sqrt(abs(sin(a[i] * b[i])));
    else
        for (int j = 0; j < KERNEL_FOR_ITERS; j++)
            c[i] += sqrt(abs(cos(a[i] * b[i])));
}

我的问题是，我能否在不影响太多速度的情况下，在整个内核执行期间防止我的屏幕变黑？

Answer 1

如果您描述您的设置，包括您正在运行的操作系统和 GPU，GPU 是否驱动显示器，以及如果在 Windows 操作系统上，GPU 处于 WDDM 或 TCC 模式，这会有些用处。

但是我们可以在没有它的情况下做出一些一般性的陈述。

正如评论中所指出的，目前，运行 CUDA 内核的 GPU 将不会为显示请求提供服务，如果它也支持显示的话。这意味着当 GPU 内核正在运行时，显示器会出现“卡住”或者可能变黑。这当然有可能在未来发生变化，但这是当前和预期的行为。

在这种情况下，如果您根本不想打扰显示器，通常的建议是使用第二个 GPU 来运行 CUDA，如果您在 Windows 上，最好是该 GPU 能够运行并放置在TCC模式。

为了在仅使用单个 GPU 时减轻影响，并确实在单 GPU 显示环境中为生产目的提供 CUDA 支持，重要的是应用程序的 CUDA 端的设计方式应使内核持续时间是有限的。为了获得良好的交互性，一个合理的起点是将内核持续时间限制为 0.1 秒或更短，因为该级别的交互性损失可能不是特别明显。如果您或某人不同意该人为因素声明，没关系；我们不必为此争论。将内核持续时间减少到您决定的任何级别都会产生良好的显示交互性。

WDDM 命令批处理在 Windows 情况下(据我所知，在 Linux 情况下不是)情况更加复杂。为了提高性能，可以对命令进行批处理，并且背靠背内核调用的批处理可能会导致交互性丢失的时间比单个内核调用所指示的时间更长。据我所知，没有任何方法可以正式解决这个问题。在每次内核调用之后，您可以通过发出伪造的(即没有必要的)CUDA 操作(例如 cudaStreamQuery())来“刷新”WDDM 命令队列。同样，我不知道正式记录的方法，并且在某种程度上它可能取决于您的应用程序设计。

关于性能，CUDA 内核启动通常涉及大约 100 微秒或更少的启动开销(我们可以称之为浪费时间)。因此，如果我们将长时间运行的内核分解为 100 毫秒的“ block ”，并且每个 block 增加约 100 微秒的开销，那么对性能的净影响可能是 CUDA 计算吞吐量减少 0.1%(假设显示任务很简单)。

以您提供的代码为例，您可能希望将该内核分解为一系列内核，在您选择的 GPU 上对其进行基准测试/计时，以便内核运行不超过大约 100 毫秒(或您选择的数量)。

#define KERNEL_FOR_ITERS 1e6


__global__ void addKernel(float *c, const float *a, const float *b,const int iters)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i % 2 == 0)
        for (int j = 0; j < iters; j++)
            c[i] += sqrt(abs(sin(a[i] * b[i])));
    else
        for (int j = 0; j < iters; j++)
            c[i] += sqrt(abs(cos(a[i] * b[i])));
}

...
const int loop_iters = 1e4; // chosen by tuning or benchmarking
cudaStream_t str;
cudaStreamCreate(&str);
for (int i = 0; i < KERNEL_FOR_ITERS; i+= loop_iters){
  addKernel<<<...,0,str>>>(d_c, d_a, d_b, loop_iters);
  cudaStreamQuery(str);//probably unnecessary on linux}

我不认为这是您实际使用的内核，但顺便说一句，可以通过将线程之间实际不同的内容限制为一小段代码来提高其性能特征。例如:

__global__ void addKernel(float *c, const float *a, const float *b,const int iters)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    float val = a[i] * b[i];
    if (i % 2 == 0)
       val = sin(val);
    else 
       val = cos(val);
    for (int j = 0; j < iters; j++)
            c[i] += sqrt(abs(val));
}

无论如何，编译器可能会计算出这种收缩，但我通常会尽量给它最好的“领先优势”。