不同 NVIDIA GPU 上具有无限循环的 cuda 内核的奇怪行为答案

【问题标题】：Strange behaviors of cuda kernel with infinite loop on different NVIDIA GPU不同 NVIDIA GPU 上具有无限循环的 cuda 内核的奇怪行为
【发布时间】：2022-01-01 23:28:07
【问题描述】：

#include <cstdio>
__global__ void loop(void) {
    int smid = -1;
    if (threadIdx.x == 0) {
        asm volatile("mov.u32 %0, %%smid;": "=r"(smid));
        printf("smid: %d\n", smid);
    }
    while (1);
}

int main() {
    loop<<<1, 32>>>();
    cudaDeviceSynchronize();
    return 0;
}

这是我的源代码，内核只是在线程索引为0时打印smid然后进入无限循环，主机只是调用之前的cuda内核并等待它。我在 2 种不同的配置下进行了一些实验，如下所示：

1。 GPU(Geforce 940M) 操作系统(Ubuntu 18.04) MPS(启用) CUDA(v11.0)
2。 GPU(Geforce RTX 3050Ti Mobile) 操作系统(Ubuntu 20.04) MPS(启用) CUDA(v11.4)

实验 1：当我在 配置 1 下运行此代码时，GUI 系统似乎被冻结，因为无法再观察到任何图形响应，但是当我按下 ctrl +c，当 CUDA 进程被杀死时，这种现象就会消失。

实验2：当我在配置2下运行这段代码时，系统似乎运行良好，没有任何异常现象，smid的输出如smid: 2\n可以显示。

实验 3：当我更改块配置 loop<<<1, 1024>>> 并在 配置 2 下运行此新代码两次时，我得到相同的 smid 输出，例如 smid: 2\nsmid: 2\n .（对于Geforce RTX 3050Ti Mobile，SM数量为20，每个多处理器的最大线程数为1536，每个块的最大线程数为1024。）

我对这些结果感到困惑，以下是我的问题：

1、为什么配置1下系统不输出smid？
2. 为什么在配置 1 下 GUI 系统似乎卡住了？
3. 与实验1不同，为什么实验2正常输出smid？
4、第三次实验，block配置达到1024个线程，这意味着不能将两个不同的block调度到同一个SM。在 MPS 环境下，所有的 CUDA 上下文将合并到一个 CUDA 上下文中，并且不再使用时间片共享 GPU 资源，但是为什么我在第三次实验中仍然得到相同的 smid？（此外，我将网格配置更改为 10 并运行它两次，smid 从 0 到 19 变化，每个 smid 只出现一次！）

【问题讨论】：

在没有无限循环的配置 1 中会发生什么。那你得到一个小输出吗？是否可以在内核调用后写入内核中的内存位置并将其复制到主机以创建内核的副作用（printf 除外）？
行为差异有多少来自驱动程序、编译器、硬件？能否为这两种架构编译或交换显卡？
感谢您的回复@Sebastian，在没有无限循环的配置 1 中，我可以获得 smid 输出。至于将smid从设备复制到主机，我稍后再试。
配置1的nvidia驱动在440左右（我记不太清楚了），配置2的驱动是495。编译器都是g++-8。也许我可以在CUDAv11.4下为sm_50构建代码，我稍后会尝试。（Geforce 940M 的计算能力为 5.0，但在 cuda11.0 中有警告说 sm_50 可能在进一步的 CUDA 版本中被弃用）

标签： c++ cuda

【解决方案1】：

为什么配置1下系统不输出smid？

一个安全的rule of thumb 是，与主机代码不同，内核中的printf 输出不会在遇到语句时打印到控制台，而是在内核和设备同步完成时与主持人。这是配置 1 中有效的实际方案，它使用的是 maxwell gpu。所以在配置 1 中没有观察到printf 输出，因为内核永远不会结束。

为什么 GUI 系统在配置 1 下似乎被冻结？

出于本次讨论的目的，有两种可能的机制：一种是前帕斯卡机制，其中compute-preemption 是不可能的，另一种是后帕斯卡机制，其中它是可能的。您的配置 1 是一个 maxwell 设备，它是 pre-pascal。您的配置 2 是安培设备，它是后帕斯卡的。所以在配置 2 中，计算抢占是有效的。这会产生多种影响，其中之一是 GPU 将“同时”满足 GUI 需求和计算内核需求（底层行为没有完整记录，而是一种时间片形式，交替关注计算内核和 GUI）。因此，在配置 1（pre-pascal）中，运行任何明显时间的内核都会在内核执行期间“冻结”GUI。在 config2 中，GPU 在某种程度上同时服务于两者。

与实验 1 不同，为什么实验 2 正常输出 smid？

虽然没有详细记录，但计算抢占过程似乎引入了一个额外的同步点，允许刷新 printf 缓冲区，如第 1 点所述。如果您阅读我在此处链接的文档，您将看到“同步点”涵盖了多种可能性，并且计算抢占似乎引入了（一种新的）一种可能性。

抱歉，目前无法回答您的第 4 个问题。 SO 的最佳实践是每个问题问一个问题。但是，我认为将 MPS 与同时为显示器提供服务的 GPU 一起使用是“不寻常的”。由于我们已经确定计算抢占在这里有效，可能是由于计算抢占以及服务显示器的需要，GPU 以循环时间片方式为客户端提供服务（因为无论如何它都必须这样做维修显示器）。在这种情况下，MPS 下的行为可能会有所不同。计算抢占允许您描述的通常限制无效。一个内核可以完全替代另一个内核。

【讨论】：