为何这段CUDA代码陷入无限循环？并发核函数未按预期执行

这段CUDA代码会无限运行，原因是kernel_loop核函数陷入无限循环。按预期，两个小型核函数应能并发启动，但实际并未达到预期效果。

#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>


int *d_x;
static constexpr int N = 1;

__global__ void init_buff(int *buff) {
  for (int i = 0; i < N; i++) {
    buff[i] = i;
  }
}

__global__ void kernel_loop(volatile int *buff) {
  while (true) {
    __threadfence();
    if (buff[0]) {
      break;
    }
  }
}

__global__ void kernel_write(volatile int *buff) {
  buff[0] = 1;
}

int main() {
  cudaMalloc(reinterpret_cast<void **>(&d_x), sizeof(int) * N);

  init_buff<<<1, 1>>>(d_x);

  cudaDeviceSynchronize();

  cudaStream_t stream1, stream2;
  cudaStreamCreateWithFlags(&stream1, cudaStreamDefault);
  cudaStreamCreateWithFlags(&stream2, cudaStreamDefault);

  cudaDeviceSynchronize();


  kernel_loop<<<1, 1, 0, stream1>>>(d_x);
  kernel_write<<<1, 1, 0, stream2>>>(d_x);

  cudaDeviceSynchronize();

  return 0;
}

测试场景与结果 #

• 调换核函数启动顺序后，程序可正常运行至结束，修改后的启动代码如下：

  kernel_write<<<1, 1, 0, stream2>>>(d_x);
  kernel_loop<<<1, 1, 0, stream1>>>(d_x);
  

• 使用cudaStreamNonBlocking标志创建流时，仍会出现无限循环，流创建代码如下：

  cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking);
  cudaStreamCreateWithFlags(&stream2, cudaStreamNonBlocking);
  

• 尝试让kernel_loop中的循环线程调用__nanosleep，无法解决该无限循环问题。

内容的提问来源于stack exchange，提问作者Elvir Crncevic