CUDA 中的全局设备变量：不好的做法？

Question

我正在设计一个包含大量 CUDA 内核的库来执行并行计算。所有内核都将作用于一个公共对象，比如一个计算网格，它是使用 C++ 样式对象定义的。计算域不一定需要从主机端访问，因此在设备端创建它并将其保留在那里目前是有意义的。我想知道是否考虑以下 "good practice"：

假设我的计算网格 class 称为 Domain。首先我定义一个全局的设备端变量来存储计算域：
```
__device__ Domain* D
```

然后我使用CUDA内核初始化计算域

__global__ void initDomain(paramType P){
    D = new Domain(P);
}

然后，我使用这个域和其他内核进行计算：

__global__ void doComputation(double *x,double *y){
    D->doThing(x,y);
    //...
}

如果我的域保持固定（即内核一旦创建就不会修改域），这样可以吗？有没有更好的办法？我最初尝试在主机端创建 Domain 对象并将其复制到设备，但结果很麻烦，因为 Domain 是一个相对复杂的类型，很难复制使用例如cudaMemCpy 甚至 Thrust::device_new（至少，我不能让它很好地工作）。

Answer 1

是的，没关系。

也许您可以使用

提高性能

__constant__

使用此关键字，您的对象将在所有内核中以非常快的内存可用。

为了复制你的对象，你必须使用：cudaMemcpyToSymbol，请注意有限制：你的对象在你的设备代码中是只读的，它必须没有默认构造函数。

您可以找到信息here

如果您的对象很复杂且难以复制，也许您可以寻找：Unified memory，然后将您的变量按值传递给内核。

CUDA 中的全局设备变量：不好的做法？

Global device variables in CUDA: bad practice?

cuda