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

Question

我的问题是在开始执行内核后，在CUDA（特别是开普勒或较新的nvidia体系结构）中调度线程块。

根据我对开普勒架构的理解（可能不正确），可以随时将单个SM调度的活动块的数量受到限制（如果我没记错的话，可以分配16个块）。同样，据我了解，一旦将块计划在特定的SM上运行，它们就无法移动。

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

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

另外，我听说块调度已“固定”到单个SM。我假设仅在块变为活动状态后才将其固定为单个SM。是这样吗

Answer 1

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

正如评论中指出的那样，如果您现在要求公开文档来支持此断言，则不确定是否可以指向它。但是，可以创建一个测试用例并向自己证明。

简而言之，您将创建一个块专用内核，该内核将启动许多块。每个SM上的第一个块将使用原子发现并声明自己。这些块将使用块已完成的计数器（再次使用原子，类似于线程防护减少示例代码）来“持久”直到所有其他块都已完成。不是首先在给定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上测试了以上代码。它没有挂起。

为了回答您的第二个问题，通常 在第一次调度的SM上保留一个块。CUDA动态并行性是一种可能的例外。