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

Posted 2023-03-06

【中文标题】在 CUDA 内核启动后，线程块调度到特定 SM 的行为是啥？

【英文标题】：What is the behavior of thread block scheduling to specific SM's after CUDA kernel launch?在 CUDA 内核启动后，线程块调度到特定 SM 的行为是什么？

【发布时间】：2015-05-20 22:28:17

【问题描述】：

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

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

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

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

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

【问题讨论】：

您所说的“SIMD”通常称为 SM（流式多处理器），或者更具体地说，对于 Kepler 来说是 SMX，对于 Maxwell 来说是 SMM。线程块彼此独立调度，并在执行资源可用时分配给 SM。此调度的确切细节是特定于实现的。你不应该依赖任何特定的行为。

“此调度的确切细节是特定于实现的。”实际调度的具体实现如何？英伟达公开声称有很多事情是特定于实现的，但私下说它们不太可能改变。

如果您确实是指 SM 而不是 SIMD，您可以编辑您的问题吗？

【参考方案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 动态并行的情况下。

【讨论】：

感谢您的代码。是否有任何方法可以完全“阻止”某些 SM（例如，可能以某种方式保持 SM 总数的一半完全被占用），以便可以使用剩余的 SM 测试应用程序？谢谢。

假设您知道您有 10 个 SM。启动一个有 20 个块的内核，每个块有 1024 个线程。使该内核代码使得如果__smid __smid 大于或等于 5，则这些块会持续 1 秒，然后退出。你最终会得到一个内核，它在 5 个 SM 上分别驻留 2 个块，充分利用这些 SM（就驻留线程或驻留线程而言，完全防止任何其他块被存放）和 5 个其他“空”的 SM .在 CUDA MPS 下执行此操作，10 个 SM 中有 5 个可用。

我已经按照您对 40GB A100 GPU 的 10GB MIG 实例的建议进行了测试；但是，***.com/questions/66668224/… 中描述了一个问题。如果您能为我提供一些帮助，那就太好了。提前致谢。