Concurrent kernel execution

[kproduct]

Hi, I’m asking help for how to correctly execute multiple kernels in different CUDA streams simultaneously.

The code I listed below seems strange to me. If I use same global function entry k() to start kernel, the whole program runs normally. However, if I change it to two different kernel entries k1() and k2(), the program hang…

#include <iostream>

__device__ volatile int s = 0;

__global__ void k1(){
  while (s == 0) {};
}

__global__ void k2(){
  s = 1;
}

__global__ void k(int x) {
  if (x == 0) {
    while (s == 0) {};
  } else {
    s = 1;
  }
}

int main() {
  cudaStream_t s1, s2;
  cudaStreamCreate(&s1);
  cudaStreamCreate(&s2);
#if 1 //!!!hang
  k1<<<1,1,0,s1>>>();
  k2<<<1,1,0,s2>>>();
#else // works
  k<<<1,1,0,s1>>>(0);
  k<<<1,1,0,s2>>>(1);
#endif
  cudaDeviceSynchronize();
  cudaError_t err = cudaGetLastError();
  if (err != cudaSuccess)
    std::cout << cudaGetErrorString(err) << std::endl;
}

The GPU used is A100 PCIe 40G with detailed info:

./deviceQuery Starting...

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

Detected 1 CUDA Capable device(s)

Device 0: "NVIDIA A100-PCIE-40GB"
  CUDA Driver Version / Runtime Version          12.2 / 12.2
  CUDA Capability Major/Minor version number:    8.0
  Total amount of global memory:                 40339 MBytes (42298834944 bytes)
  (108) Multiprocessors, (064) CUDA Cores/MP:    6912 CUDA Cores
  GPU Max Clock rate:                            1410 MHz (1.41 GHz)
  Memory Clock rate:                             1215 Mhz
  Memory Bus Width:                              5120-bit
  L2 Cache Size:                                 41943040 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total shared memory per multiprocessor:        167936 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 3 copy engine(s)
  Run time limit on kernels:                     No
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Enabled
  Device supports Unified Addressing (UVA):      Yes
  Device supports Managed Memory:                Yes
  Device supports Compute Preemption:            Yes
  Supports Cooperative Kernel Launch:            Yes
  Supports MultiDevice Co-op Kernel Launch:      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 225 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 12.2, CUDA Runtime Version = 12.2, NumDevs = 1
Result = PASS

I did a lot experiments to try concurrent kernel execution, and finally found this modification matters. Could you give me some advice? Thanks!

[striker159]

First of all, there is no guarantee by CUDA that two kernels in different streams will execute concurrently.

Your specific problem is probably caused by lazy loading. Simply speaking, k1 needs to complete before k2 can be loaded, but k1 cannot complete unless k2 is complete, which creates a deadlock.
Try setting the environment variable CUDA_MODULE_LOADING to EAGER

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#concurrent-execution

Loading kernels might require context synchronization. Some programs incorrectly treat the possibility of concurrent execution of kernels as a guarantee. In such cases, if program assumes that two kernels will be able to execute concurrently, and one of the kernels will not return without the other kernel executing, there is a possibility of a deadlock.