gpgpu - 关于CUDA从block到SM的分布细节问题

Question

以计算能力为1.3的硬件为例。

30 条短信可用。那么最多可以同时运行240个块(考虑到寄存器和共享内存的限制，对块数的限制可能会低很多)。超过 240 的那些块必须等待可用的硬件资源。

我的问题是那些超过 240 的块何时会分配给 SM。曾经一些 前 240 块已完成？或者当 全部 前 240 个块完成了吗？

我写了这么一段代码。

#include<stdio.h>
#include<string.h>
#include<cuda_runtime.h>
#include<cutil_inline.h>

const int BLOCKNUM = 1024;
const int N=240;
__global__ void kernel ( volatile int* mark ) {
    if ( blockIdx.x == 0 ) while ( mark[N] == 0 );
    if ( threadIdx.x == 0 ) mark[blockIdx.x] = 1;
}

int main() {
    int * mark;
    cudaMalloc ( ( void** ) &mark, sizeof ( int ) *BLOCKNUM );
    cudaMemset ( mark, 0, sizeof ( int ) *BLOCKNUM );
    kernel <<< BLOCKNUM, 1>>> ( mark );
    cudaFree ( mark );
    return 0;
}

此代码会导致死锁并无法终止。但是如果我将 N 从 240 更改为 239，代码就可以终止。所以我想知道一些关于块调度的细节。

Answer 1

在 GT200 上，已经通过微基准测试证明，只要 SM 退出了它正在运行的所有当前事件块，就会调度新块。所以答案是当一些块完成时，调度粒度是SM级别的。似乎有一个共识，即 Fermi GPU 比前几代硬件具有更精细的调度粒度。