WSL2中cudaMemPrefetchAsync位置引发CUDA Error 101的原因分析

问题描述 #

在WSL环境下使用CUDA 12.1运行两段仅cudaMemPrefetchAsync调用位置不同的代码：

• snip1中cudaMemPrefetchAsync在核函数启动前调用，返回CUDA error 101: invalid device ordinal错误
• snip2中cudaMemPrefetchAsync在核函数启动后调用，可正常运行

请问该差异是否完全由函数位置导致？背后的机制是什么？

snip1（报错代码） #

#include <curand_kernel.h>

#include <cstdio>

void checkError()
{
  cudaError_t err_;
  err_ = cudaGetLastError();
  if (err_ != cudaSuccess)
  {
    std::printf("CUDA error %d:%s at %s:%d\n", err_, cudaGetErrorString(err_), __FILE__, __LINE__);
    exit(EXIT_FAILURE);
  }
}

__global__ void init_curand(curandState *states, unsigned long long seed)
{
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  int j = threadIdx.y + blockIdx.y * blockDim.y;
  int idx = i * blockDim.y * gridDim.y + j;
  curand_init(seed, idx, 0, &states[idx]);
}

int main()
{
  int deviceId;
  int numberOfSMs;

  cudaGetDevice(&deviceId);
  cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);
  printf("Device ID: %d\tNumber of SMs: %d\n", deviceId, numberOfSMs);

  dim3 threadsPerBlock(16, 16);
  dim3 numBlocks(8 * numberOfSMs, 8 * numberOfSMs);

  const int M = 10;
  const int N = 20;
  const int bytes = M * N * sizeof(int8_t);
  int8_t *noisy;
  int8_t *ising1;
  int8_t *ising2;
  cudaMallocManaged(&noisy, bytes);
  cudaMallocManaged(&ising1, bytes);
  cudaMallocManaged(&ising2, bytes);
  curandState *states;
  cudaMalloc(&states, numBlocks.x * threadsPerBlock.x * numBlocks.y * threadsPerBlock.y *
                          sizeof(curandState));

  /* ??? */
  cudaMemPrefetchAsync(noisy, bytes, deviceId);
  cudaMemPrefetchAsync(ising1, bytes, deviceId);
  cudaMemPrefetchAsync(ising2, bytes, deviceId);

  init_curand<<<numBlocks, threadsPerBlock>>>(states, time(NULL));
  checkError();

  return 0;
}

snip2（正常代码） #

#include <curand_kernel.h>

#include <cstdio>

void checkError()
{
  cudaError_t err_;
  err_ = cudaGetLastError();
  if (err_ != cudaSuccess)
  {
    std::printf("CUDA error %d:%s at %s:%d\n", err_, cudaGetErrorString(err_), __FILE__, __LINE__);
    exit(EXIT_FAILURE);
  }
}

__global__ void init_curand(curandState *states, unsigned long long seed)
{
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  int j = threadIdx.y + blockIdx.y * blockDim.y;
  int idx = i * blockDim.y * gridDim.y + j;
  curand_init(seed, idx, 0, &states[idx]);
}

int main()
{
  int deviceId;
  int numberOfSMs;

  cudaGetDevice(&deviceId);
  cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);
  printf("Device ID: %d\tNumber of SMs: %d\n", deviceId, numberOfSMs);

  dim3 threadsPerBlock(16, 16);
  dim3 numBlocks(8 * numberOfSMs, 8 * numberOfSMs);

  const int M = 10;
  const int N = 20;
  const int bytes = M * N * sizeof(int8_t);
  int8_t *noisy;
  int8_t *ising1;
  int8_t *ising2;
  cudaMallocManaged(&noisy, bytes);
  cudaMallocManaged(&ising1, bytes);
  cudaMallocManaged(&ising2, bytes);

  curandState *states;
  cudaMalloc(&states, numBlocks.x * threadsPerBlock.x * numBlocks.y * threadsPerBlock.y *
                          sizeof(curandState));

  init_curand<<<numBlocks, threadsPerBlock>>>(states, time(NULL));
  checkError();

  /* ??? */
  cudaMemPrefetchAsync(noisy, bytes, deviceId);
  cudaMemPrefetchAsync(ising1, bytes, deviceId);
  cudaMemPrefetchAsync(ising2, bytes, deviceId);

  return 0;
}

解答 #

结论 #

是的，错误完全由cudaMemPrefetchAsync的调用位置差异导致。

背后机制 #

1. 设备上下文初始化时机：cudaGetDevice仅获取当前默认设备的ID，但不会主动初始化该设备的运行上下文。设备上下文的初始化通常由首次执行核函数、显式同步操作或某些设备内存操作触发。
2. snip1的错误原因：在核函数启动前调用cudaMemPrefetchAsync时，设备上下文尚未初始化。在WSL环境的CUDA 12.1实现中，异步预取操作需要关联到已初始化的设备上下文，此时驱动会因上下文未就绪，错误判定传入的deviceId无效，返回invalid device ordinal错误。
3. snip2的正常原因：先启动核函数init_curand会强制触发设备上下文的初始化，此时设备已处于活跃状态。后续调用cudaMemPrefetchAsync时，deviceId能正确关联到已初始化的设备上下文，异步预取操作可正常执行。
4. 异步操作的依赖：cudaMemPrefetchAsync是异步操作，依赖CUDA流（默认是流0）调度执行。如果设备上下文未初始化，流无法绑定到有效设备，进而引发设备序数无效的错误。

内容的提问来源于stack exchange，提问作者CHEN Yunkai