PyCUDA 中的内存分配顺序是否重要？

Question

我在 CUDA 统一内存 [2] 上使用 PyCUDA 的接口 [1]。在某些时候，我添加了随机数生成器 [3] 并盯着 Jupyter Notebook 中的死内核：

我将问题缩小到创建随机数生成器。或者，准确地说，到我这样做的那一刻：

import pycuda.curandom
from pycuda import autoinit, driver
import numpy as np

gpu_data_1 = driver.managed_zeros(shape=5, dtype=np.int32, mem_flags=driver.mem_attach_flags.GLOBAL)
gpu_generator = pycuda.curandom.XORWOWRandomNumberGenerator(pycuda.curandom.seed_getter_uniform)
gpu_data_2 = driver.managed_zeros(shape=5, dtype=np.int32, mem_flags=driver.mem_attach_flags.GLOBAL)

上面的代码在没有任何错误消息的情况下失败，但如果我将gpu_generator = ... 行高或低一行，它似乎工作正常。

我相信 PyCUDA 可能无法执行 prepare call，这归结为这个内核：

extern "C" {
    __global__ void prepare(curandStateXORWOW *s, const int n,
        unsigned int *v, const unsigned int o)
    {
      const int id = blockIdx.x*blockDim.x+threadIdx.x;
      if (id < n)
        curand_init(v[id], id, o, &s[id]);
    }
}

知道可能是什么问题吗？

Answer 1

在前 Pascal UM（统一内存）机制中，主机代码在内核启动后、cudaDeviceSynchronize() 发出之前接触托管分配是 illegal。

我猜这个代码违反了这个规则。如果我在 Maxwell 系统上运行您的复制案例，我会得到：

$ cuda-memcheck python ./idontthinkso.py
========= CUDA-MEMCHECK
========= Error: process didn't terminate successfully
========= Fatal UVM CPU fault due to invalid operation
=========     during write access to address 0x703bc1000
=========
========= ERROR SUMMARY: 1 error

那是托管内存系统崩溃了。在随机生成器设置（运行内核）和 zeros 调用（涉及托管内存）之间进行同步调用可以在我的系统上消除它：

$ cat idontthinkso.py 
import pycuda.curandom
from pycuda import autoinit, driver
import numpy as np

gpu_data_1 = driver.managed_zeros(shape=5, dtype=np.int32, mem_flags=driver.mem_attach_flags.GLOBAL)
gpu_generator = pycuda.curandom.XORWOWRandomNumberGenerator(pycuda.curandom.seed_getter_uniform)
autoinit.context.synchronize()
gpu_data_2 = driver.managed_zeros(shape=5, dtype=np.int32, mem_flags=driver.mem_attach_flags.GLOBAL)

$ cuda-memcheck python ./idontthinkso.py
========= CUDA-MEMCHECK
========= ERROR SUMMARY: 0 errors

您所处的 UM 制度会因您使用的 GPU、驱动程序和操作系统而异。