非常驻线程块的共享内存在哪里？

Question

我正在尝试了解共享内存的工作原理，当块大量使用它时。

所以我的 gpu (RTX 2080 ti) 每个 SM 有 48 kb 的共享内存，每个线程块也是如此。在下面的示例中，我在同一个 SM 上强制使用了 2 个块，每个块使用完整的 48 kb 内存。我强制两个块在完成之前进行通信，但由于它们不能并行运行，这应该是一个死锁。然而，无论我运行 2 个块还是 1000 个块，程序都会终止。

这是因为块 1 一旦遇到死锁就暂停，并与块 2 切换？如果是，那么当块 2 处于活动状态时，来自块 1 的 48 kb 数据去哪里了？是否存储在全局内存中？

内核：

__global__ void testKernel(uint8_t* globalmem_message_buffer, int n) {
    const uint32_t size = 48000;
    __shared__ uint8_t data[size];
    for (int i = 0; i < size; i++) 
        data[i] = 1;

    globalmem_message_buffer[blockIdx.x] = 1;
    while (globalmem_message_buffer[(blockIdx.x + 1) % n] == 0) {}
    printf("ID: %d\n", blockIdx.x);

}

主机代码：

    int n = 2; // Still works with n=1000
    cudaStream_t astream;
    cudaStreamCreate(&astream);
    uint8_t* globalmem_message_buffer;
    cudaMallocManaged(&globalmem_message_buffer, sizeof(uint8_t) * n);
    for (int i = 0; i < n; i++) globalmem_message_buffer[i] = 0;
    cudaDeviceSynchronize();
    testKernel << <n, 1, 0, astream >> > (globalmem_message_buffer, n);

编辑：将“threadIdx”更改为“blockIdx”

Answer 1

所以我的 gpu (RTX 2080 ti) 每个 SM 有 48 kb 的共享内存，每个线程块也是如此。在下面的示例中，我在同一个 SM 上强制使用了 2 个块，每个块使用完整的 48 kb 内存。

那不会发生。这里的一般前提是有缺陷的。当有足够的空闲资源支持该块时，GPU 块调度程序仅在 SM 上放置一个块。

具有 48KB 共享内存的 SM 上已经有一个使用 48KB 共享内存的块驻留，在现有/驻留块“退休”并释放之前，不会将任何该类型的新块存放在其上它正在使用的资源。

因此，在正常的 CUDA 调度模型中，一个块可以非常驻的唯一方法是它从未在 SM 上被调度过。在这种情况下，它在队列中等待时不使用任何资源。

在 CUDA 抢占的情况下例外。这种机制没有很好的文档记录，但会在例如上下文切换时发生。在这种情况下，整个线程块状态以某种方式从 SM 中删除并存储在其他地方。然而，抢占不适用于我们分析单个内核启动行为的情况。

您尚未提供完整的代码示例，但是，对于 n=2 案例，您声称这些代码会以某种方式存放在同一个 SM 上的说法根本不正确。

对于n=1000 的情况，您的代码只需要将内存中的单个位置设置为1：

while (globalmem_message_buffer[(threadIdx.x + 1) % n] == 0) {}

threadIdx.x 因为您的代码始终为 0，因为您正在启动只有 1 个线程的线程块：

testKernel << <n, 1, 0, astream >> > (globalmem_message_buffer, n);

因此这里生成的索引总是1（对于n大于等于2）。所有线程块都在检查位置 1。因此，当 blockIdx.x 为 1 的线程块执行时，网格中的所有线程块都将“解锁”，因为它们都在测试同一个位置。简而言之，您的代码可能没有按照您的想法或预期进行。即使您让每个线程块检查另一个线程块的位置，我们也可以想象一系列线程块存款可以满足这一点，而无需同时驻留所有n 线程块，所以我认为这也不能证明任何事情。 （区块存款顺序没有指定顺序。）