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

Question

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

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

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

是否在 SM 中一个当前运行ning 活动块完成后立即执行新块？还是仅在 SM 完成其所有当前活动块后才执行下一组块？还是只有在所有 SM 完成所有当前活动的块执行后它们才开始？

另外我听说块调度是"fixed"到单个SM。我假设它仅在块激活后才固定到单个 SM。是这样吗？

Answer 1

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

正如评论中所指出的，如果您现在要求 public 文档来支持此断言，我不确定我是否可以指出它。然而，可以创建一个测试用例并向自己证明这一点。

简而言之，您将创建一个将启动许多块的块专用内核。每个 SM 上的第一个块将使用原子发现并声明自己。这些块将 "persist" 直到所有其他块都完成，使用块完成计数器（同样，使用原子，类似于 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 on the SM on which it was first scheduled. One possible exception是在CUDA动态并行的情况下

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

What is the behavior of thread block scheduling to specific SM's after CUDA kernel launch?

cuda