多个cuda内核可以在同一个SM上并行执行吗？

Question

假设一个 cuda gpu 只有一个 SM。

我可以并行启动两个使用 512 个线程的 cuda 内核吗？或者内核是按块分配的，导致内核串行执行而不是并行执行，从而使剩余的 512 个内核闲置以进行两次执行？

我尝试查找描述 nvidia gpus 背后的调度机制的文档，但我只能找到描述单个内核如何获取分配给它的资源的文档。它没有说明是否可以将两个内核并行分配给单个块，或者是否在扭曲级别甚至线程级别分配资源。

我期望块级或扭曲级资源共享，但我想了解实际的调度算法，而不是猜测或逆向工程。

Answer 1

我稍微编辑了这个问题，使其或多或少合理，并且不会将软件层次结构（块）与硬件层次结构（SM、cuda 核心等）混为一谈。在我看来，评论流澄清了 OP 的确认.

是的，即使在单个 SM 上，块调度程序也可以从两个不同的内核存储块。

这是一个示例/测试用例：

$ /usr/local/cuda/samples/bin/x86_64/linux/release/deviceQuery
/usr/local/cuda/samples/bin/x86_64/linux/release/deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "Quadro K610M"
  CUDA Driver Version / Runtime Version          11.4 / 11.4
  CUDA Capability Major/Minor version number:    3.5
  Total amount of global memory:                 981 MBytes (1028784128 bytes)
  (001) Multiprocessors, (192) CUDA Cores/MP:    192 CUDA Cores
  GPU Max Clock rate:                            954 MHz (0.95 GHz)
  Memory Clock rate:                             1300 Mhz
  Memory Bus Width:                              64-bit
  L2 Cache Size:                                 524288 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
  Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 16384), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total shared memory per multiprocessor:        49152 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  2048
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 1 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device supports Managed Memory:                Yes
  Device supports Compute Preemption:            No
  Supports Cooperative Kernel Launch:            No
  Supports MultiDevice Co-op Kernel Launch:      No
  Device PCI Domain ID / Bus ID / location ID:   0 / 1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 11.4, CUDA Runtime Version = 11.4, NumDevs = 1
Result = PASS
$ cat t37.cu
#include <iostream>
__device__ volatile int s = 0;
__global__ void k1(){
  while (s == 0) {};
}
__global__ void k2(){
  s = 1;
}

int main(){

  cudaStream_t s1, s2;
  cudaStreamCreate(&s1);
  cudaStreamCreate(&s2);
  k1<<<1,1,0,s1>>>();
  k2<<<1,1,0,s2>>>();
  cudaDeviceSynchronize();
  cudaError_t err = cudaGetLastError();
  if (err != cudaSuccess) std::cout << cudaGetErrorString(err) << std::endl;
}
$ nvcc -o t37 t37.cu -arch=sm_35
nvcc warning : The 'compute_35', 'compute_37', 'compute_50', 'sm_35', 'sm_37' and 'sm_50' architectures are deprecated, and may be removed in a future release (Use -Wno-deprecated-gpu-targets to suppress warning).
$ ./t37
$

这是在配备 Quadro K610M 的笔记本电脑上。

deviceQuery

向我们展示了该 GPU 有一个 SM。

测试代码显示了两个内核。第一个启动的内核 (

k1

) 将永远旋转，直到全局位置 (

s

) 变为非零。第二个启动的内核 (

k2

) 将全局位置设置为非零，允许第一个内核退出。最终结果是应用程序正常终止。如果两个内核没有在同一个 SM 上同时调度（只有 1 个），那么第一个内核将永远旋转，第二个内核将永远不会运行，并且观察到的应用程序行为将是挂起。 （例如，您可以通过将两个内核启动到同一流（例如默认的 NULL 流）来见证挂起。）

我们没有看到挂起的事实意味着两个内核同时在同一个 SM 上运行。

SO 标签上还有许多其他问题，讨论块调度和 GPU 调度行为。

这里

是一个例子。还有其他的。