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

Question

我的问题是在内核已经开始执行之后，CUDA（特别是 kepler 或更新的 nvidia 架构）中的线程块调度。

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

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

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

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

Answer 1

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

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

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

这是一个完整的例子：

$ 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 测试了上述代码。它不会挂起。

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