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

Question

我正在设计一个库，其中包含大量 CUDA 内核来执行并行计算。所有内核都将作用于一个公共对象，例如计算网格，它是使用 C++ 样式对象定义的。计算域不一定需要从主机端访问，因此在设备端创建它并保持它现在是有意义的。我想知道以下是否被认为是“良好做法”：

• 假设我的计算网格类名为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，然后将变量按值传递给内核。