非驻留线程块的共享内存去哪里了？

Question

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

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

是不是block 1 运行进入死锁后就暂停了，和block 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

So my gpu (RTX 2080 ti) has 48 kb of shared memory per SM, and the same per threadblock. In my example below i have 2 blocks forced on the same SM, each using the full 48 kb of memory.

那不会发生。这里的一般前提是有缺陷的。 GPU 块调度程序仅在有足够的可用资源来支持该块时才将块存放在 SM 上。

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

因此在正常的 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）。所有threadblocks都在检查location 1。因此，当blockIdx.x为1的threadblock执行时，grid中的所有threadblocks都会“unblocked”，因为它们都在测试同一个地点。简而言之，您的代码可能没有按照您的想法或预期进行。即使你让每个线程块检查另一个线程块的位置，我们可以想象一系列线程块存款可以满足这一点，而不需要所有 n 个线程块同时驻留，所以我认为这也不能证明什么. （区块充值顺序没有指定顺序。）

非驻留线程块的共享内存去哪里了？

Where do shared memory of non-resident threadblocks go?

cuda