非驻留线程块的shared memory存储位置及CUDA调度死锁问题咨询

问题核心原因拆解 #

• 第一个认知误区：你并没有实现「强制2个线程块调度到同一个SM」的效果
  RTX 2080 Ti（图灵架构）共有68个流多处理器（SM），CUDA默认调度策略会优先把不同线程块分配到空闲SM上执行，你启动2个线程块时，两个块会直接跑在两个不同的SM上，每个SM单独分配48KB共享内存给对应块使用，完全不存在资源冲突。两个块并行执行时，都会先把自己的global标记位置1，再读取对方的标记位时已经是1，直接退出循环，自然不会死锁。
• 第二个认知误区：CUDA硬件目前不支持线程块级别的上下文切换
  现有NVIDIA消费级GPU的调度逻辑是：线程块一旦被调度到某个SM上运行，就会独占它申请的所有SM资源（共享内存、寄存器等）直到执行完毕，不会被暂停切换给同SM的其他块。你假设的「块1被暂停、切换到块2运行」的情况根本不会发生。

如果你真的强制两个块绑定到同一个SM会发生什么？ #

你的每个线程块申请了48KB共享内存，刚好占满单个SM的全部共享内存容量，因此硬件调度器会判定该SM无法同时容纳两个块，只会先调度其中一个块到SM上运行，另一个块留在全局调度队列中等待。
此时先运行的块会把自己的global标记位置1，随后进入while循环等待另一个块的标记位，但另一个块永远不会被调度（当前SM资源被占满，没有空闲资源分配给它），这时才会触发真正的死锁，程序会彻底卡住。如果要复现你预期的死锁场景，可以使用CUDA 11及以上版本提供的SM亲和性接口，强制两个线程块绑定到同一个SM执行，此时程序会直接卡住。

为什么n=1000时程序仍然正常运行？ #

RTX 2080 Ti共有68个SM，每个SM一次只能运行1个你写的线程块（因为共享内存占满了），所以第一批会先调度68个块并行运行：

1. 前面的67个块（块0~块66）等待的下一个块都在第一批运行队列里，对应的标记位很快会被置1，因此这67个块执行完后会释放SM资源
2. 调度器会接着把后续的块（块68及之后）调度到空出来的SM上运行，块67等待的块68标记位被置1后也会退出
3. 以此类推，所有块会分批完成执行，最后一个块等待的块0的标记位早就被置1，因此不会出现死锁。

关于你提到的共享内存上下文存储问题 #

现有硬件不支持块级切换，因此不存在这个场景。如果是未来支持块上下文切换的专业级GPU，共享内存的上下文数据确实会被转储到全局内存或者专门的上下文保留区域，等到块被切换回SM时再恢复到片上共享内存中。

测试用内核代码 #

__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; // 改为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);

内容的提问来源于stack exchange，提问作者Daniel Johansen