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

Question

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

cuda

4

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

根据我对Kepler架构的理解（可能不正确），在任何时刻，单个SM可以调度的活动块数量存在限制（如果我记得正确，则为16个块）。此外，我了解到一旦块被安排在特定的SM上运行，它们就无法移动。

我想知道的是，在初始选择块并在设备上开始执行后，块的调度和执行行为如何（假设内核具有的线程块数大于所有SM中的活动线程块数）。

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

此外，我听说块调度被“固定”到单个SM。我假设块在激活后只固定在单个SM上。这是真的吗？

- NothingMore

1

你所提到的“SIMD”通常被称为SM（流式多处理器），或更具体地说，Kepler的SMX和Maxwell的SMM。线程块是独立调度的，并且在可用资源的情况下分配给SM以进行执行。这种调度的确切细节是特定于实现的。您不应依赖任何特定的行为。 - void_ptr

这个调度的确切细节取决于具体实现。实际调度有多少是与具体实现相关的？Nvidia公开声称有很多东西是与具体实现相关的，但在私下里却表示它们不太可能改变。 - NothingMore

2

如果您确实是指SM而不是SIMD，您可以编辑您的问题以反映这一点吗？ - Jez

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Robert Crovella · Accepted Answer

只要SM有足够的未使用资源来支持新块，就可以立即安排新块。在安排新块之前，不必将SM完全排干块。正如评论中指出的那样，如果您现在要求公共文档支持这个断言，我不确定我能指出它。但是，您可以创建一个块专用内核，启动许多块。每个SM上的第一个块将使用原子操作发现和声明自己。这些块将“持久存在”，直到所有其他块完成，使用块完成计数器（再次使用原子操作，类似于threadfence reduction示例代码）。在给定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动态并行性的一个可能的例外情况。