研究一下 GPU 线程块到 SM 的调度机制

post on 09 Mar 2025 about 1543words require 6min
CC BY 4.0 （除特别声明或转载文章外）
如果这篇文章帮助到你，可以请我喝一杯咖啡~

发现自己的知识有一些盲区，浅研究一下 GPU 线程块到 SM 的调度机制。据我所知，没有任何方式能够调度指定线程块到指定 SM 上，也没有办法能确保线程块占满所有 SM。

简单实验

// nvcc -arch=sm_80 -run a.cu
#include <cstdint>
#include <cstdio>
#include <cuda_runtime.h>
__global__ void kernel(uint32_t id) {
  __nanosleep(1e9);
  // sleep to ensure concurrent kernel execution
  // `__nanosleep` only support sm_70 and higher
  uint32_t sm_id;
  asm volatile("mov.u32 %0, %smid;" : "=r"(sm_id));
  printf("hello from kernel %u, block %u, sm %u\n", id, blockIdx.x, sm_id);
}
int main() {
  const int N = 5;
  cudaStream_t s[N];
  for (int i = 0; i < N; ++i)
    cudaStreamCreateWithFlags(&s[i], cudaStreamNonBlocking);
  for (int i = 0; i < N; ++i)
    kernel<<<1, 1, 0, s[i]>>>(i);
  for (int i = 0; i < N; ++i)
    cudaStreamSynchronize(s[i]);
  kernel<<<N, 1, 0, s[0]>>>(N);
  cudaStreamSynchronize(s[0]);
  for (int i = 0; i < N; ++i)
    cudaStreamDestroy(s[i]);
}

在 A100-PCIE-40GB, CUDA 12.6.2, driver 560.35.03 上，程序输出如下。

hello from kernel 4, block 0, sm 8
hello from kernel 0, block 0, sm 0
hello from kernel 3, block 0, sm 6
hello from kernel 1, block 0, sm 2
hello from kernel 2, block 0, sm 4
hello from kernel 5, block 0, sm 0
hello from kernel 5, block 4, sm 8
hello from kernel 5, block 3, sm 6
hello from kernel 5, block 1, sm 2
hello from kernel 5, block 2, sm 4

结论：不管是同一个 stream 的同一个 kernel，还是不同 stream 上的不同 kernel，都会优先调度到空闲的 SM 上（推测为 Round-Robin 轮询），然后才是已有的 SM 上。在分配线程块到 SM 之前，线程块调度器会进行一个占用测试，检查每个 SM 当前的资源利用情况（线程/warps 数量，寄存器，共享内存），以确定是否可以容纳新的线程块。此测试的目的是确保当前的占用率能够满足新内核的需求，从而实现线程块到 SM 的映射。此外，多次运行发现用到的 SM 都只有偶数，不知道这里有什么窍门。

研究一下 GPU 线程块到 SM 的调度机制

简单实验

相关文献

Related posts