cuda - 在 CUDA 内核启动后，线程 block 调度到特定 SM 的行为是什么？

Question

我的问题是关于在内核执行已经开始后 CUDA(特别是开普勒或更新的 nvidia 架构)中的线程 block 的调度。

根据我对开普勒架构的理解(这可能是不正确的)，可以在任何时间安排到单个 SM 的事件 block 的数量是有限的(如果我没记错的话，是 16 个 block )。同样据我了解，一旦计划在特定 SM 上运行， block 就无法移动。

我很好奇的是在 block 的初始选择发生并开始在设备上执行之后的 block 调度和执行行为(假设内核的线程 block 比所有 SM 中的活跃线程 block 多)。

一个当前正在运行的事件 block 在 SM 中完成后是否立即执行新 block ？还是只有在 SM 完成所有当前事件的 block 后才执行下一组 block ？还是仅在所有 SM 完成所有当前事件 block 执行后才启动？

此外，我听说 block 调度是“固定”到单个 SM 的。我假设它仅在 block 激活后才固定到单个 SM。是这样吗？

Answer 1

只要 SM 有足够的未使用资源来支持新 block ，就可以调度新 block 。在调度新 block 之前，没有必要让 SM 完全耗尽 block 。

正如评论中所指出的，如果您现在要求提供公共(public)文件来支持这一断言，我不确定我能否指出这一点。但是，可以创建一个测试用例并向自己证明这一点。

简而言之，您将创建一个可以启动许多 block 的 block 专用内核。每个 SM 上的第一个 block 将使用原子发现并声明自己。这些 block 将“持续”直到所有其他 block 都完成，使用 block 完成计数器(同样，使用原子，类似于 threadfence 减少示例代码)。不是第一个在给定 SM 上启动的所有其他 block 将简单地退出。这样的代码的完成，而不是挂起，将证明即使某些 block 仍然存在，其他 block 也可以被调度。

这是一个完整的示例:

$ cat t743.cu
#include <stdio.h>
#include <stdint.h>
#include <stdlib.h>

#define NB 1000
// increase array length here if your GPU has more than 32 SMs
#define MAX_SM 32
// set HANG_TEST to 1 to demonstrate a hang for test purposes
#define HANG_TEST 0

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

static __device__ __inline__ uint32_t __smid(){
    uint32_t smid;
    asm volatile("mov.u32 %0, %%smid;" : "=r"(smid));
    return smid;}

__device__ volatile int blocks_completed = 0;
// increase array length here if your GPU has more than 32 SMs
__device__ int first_SM[MAX_SM];

// launch with one thread per block only
__global__ void tkernel(int num_blocks, int num_SMs){

  int my_SM = __smid();
  int im_not_first = atomicCAS(first_SM+my_SM, 0, 1);
  if (!im_not_first){
    while (blocks_completed < (num_blocks-num_SMs+HANG_TEST));
  }
  atomicAdd((int *)&blocks_completed, 1);
}

int main(int argc, char *argv[]){
  unsigned my_dev = 0;
  if (argc > 1) my_dev = atoi(argv[1]);
  cudaSetDevice(my_dev);
  cudaCheckErrors("invalid CUDA device");
  int tot_SM = 0;
  cudaDeviceGetAttribute(&tot_SM, cudaDevAttrMultiProcessorCount, my_dev);
  cudaCheckErrors("CUDA error");
  if (tot_SM > MAX_SM) {printf("program configuration error\n"); return 1;}
  printf("running on device %d, with %d SMs\n", my_dev, tot_SM);
  int temp[MAX_SM];
  for (int i = 0; i < MAX_SM; i++) temp[i] = 0;
  cudaMemcpyToSymbol(first_SM, temp, MAX_SM*sizeof(int));
  cudaCheckErrors("cudaMemcpyToSymbol fail");
  tkernel<<<NB, 1>>>(NB, tot_SM);
  cudaDeviceSynchronize();
  cudaCheckErrors("kernel error");
}

$ nvcc -o t743 t743.cu
$ ./t743 0
running on device 0, with 15 SMs
$ ./t743 1
running on device 1, with 1 SMs
$ ./t743 2

我已经在 Linux 上使用 CUDA 7、K40c、C2075 和 Quadro NVS 310 GPU 测试了上述代码。它不挂。

回答你的第二个问题，一个 block 一般 remains在第一次安排它的 SM 上。一种可能的 exception是在 CUDA 动态并行的情况下。