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

Question

#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) OS(Ubuntu 18.04) MPS(启用) CUDA(v11.0)
2。 GPU（Geforce RTX 3050Ti 移动版）OS（Ubuntu 20.04）MPS（启用）CUDA(v11.4)

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

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

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

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

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

Answer 1

Why doesn't the system output smid under configuration 1?

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

Why does the GUI system seems to get freezed under configuration 1?

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

Unlike experiment 1, why does experiment 2 output smid normally?

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

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

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

Strange behaviors of cuda kernel with infinite loop on different NVIDIA GPU

c++

cuda