在设备代码中创建对象

Question

我想在设备上创建一个对象并将其分配给主机上可用的指针。我在这里做错了什么吗？

__global__ void createAProduction(DeviceProduction* production) {
    production = new AProduction();
}

DeviceProduction * devAProduction = NULL;
cudaMalloc(&devAProduction, sizeof(AProduction));
createAProduction<<<1, 1>>>(devAProduction);
deviceProductions["A"] = devAProduction;

在代码中的某个地方我想做某事。喜欢：

BatchOperation ** devBatchOperations;
    cudaMalloc((void **) &devBatchOperations, sizeof(BatchOperation *) * operationCount);

然后我用那个填充那个指针数组：

void DeviceBatchExecutor::execute(vector<BatchOperation> operationsToPerform) {
    BatchOperation ** devBatchOperations;
    cudaMalloc((void **) &devBatchOperations, sizeof(BatchOperation *) * operationsToPerform.size());
    int i = 0;
    for(batchOperationIt it = operationsToPerform.begin(); it != operationsToPerform.end(); ++it) {
        BatchOperation * devBatchOperation;
        cudaMalloc(&devBatchOperation, sizeof(BatchOperation));
        cudaMemcpy(&devBatchOperation, &it, sizeof(BatchOperation), cudaMemcpyHostToDevice);
        Vertex * devInputNode = it->inputNode->allocateToDevice();
        cudaMemcpy(&(devBatchOperation->inputNode), &devInputNode, sizeof(Vertex *), cudaMemcpyDeviceToDevice);
        cudaMemcpy(&(devBatchOperation->production), &(it->production), sizeof(Production *), cudaMemcpyDeviceToDevice);
        cudaMemcpy(&devBatchOperations[i], &devBatchOperation, sizeof(BatchOperation *), cudaMemcpyDeviceToDevice);
        i++;
    }
    int operationCount = operationsToPerform.size();
    executeOperations<<<operationCount, 1>>>(devBatchOperations);
}

其中 production 是指向保存创建对象 AProduction 的设备内存的指针。然后我终于通过

调用处理

executeOperations<<<operationCount, 1>>>(devBatchOperations);

所以我依赖于虚拟方法调用。由于这些 DeviceProduction 对象是在设备上创建的，因此还有一个虚拟指针表，因此它应该可以工作。请参阅示例here。但它不是，因为收到的批处理操作似乎是随机的......调用时崩溃。

__global__ void executeOperations(BatchOperation ** operation) {    
    operation[blockIdx.x]->production->apply(operation[blockIdx.x]->inputNode);
}

批处理操作是一个包含要执行的生产的结构。

struct BatchOperation {
    Production * production;
    Vertex * inputNode;
    Vertex * outputNode;
};

Answer 1

我在这里做错了吗？

是的，可能。指针production 被传递给内核按值：

createAProduction<<<1, 1>>>(devAProduction);

它指向设备内存中的某个位置，因为您已经在其上运行了cudaMalloc。这行内核代码：

production = new AProduction();

用一个新的指针覆盖production 的值传递副本，由内核中的new 返回。这几乎肯定不是你的本意。 （而且您还没有定义 AProduction 是什么。）。在该内核调用完成时，指针的按值传递“副本”无论如何都会丢失。您也许可以像这样修复它：

*production = *(new DeviceProduction());

现在您的production 指针指向设备内存中的一个区域，该区域包含一个实例化（在设备上）对象，这似乎是您的意图。创建一个新对象只是为了复制它可能没有必要，但这不是我在这里要指出的问题的症结所在。您也可以通过将指针传递给内核来“解决”这个问题。然后，您需要分配一个指针数组，并使用内核中的new 直接分配一个单独的指针，如您所示。

您的代码的其余部分有很多未定义的项目。例如，在上面的代码中，不清楚为什么要声明 production 是指向 DeviceProduction 类型的指针，然后尝试为其分配 AProduction 类型。大概这是某种尚不清楚的对象继承形式。

由于您还没有真正提供任何接近完整代码的东西，我从here 中借用了一些片段来组合一个完整的工作示例，显示在一个内核中创建/设置对象，然后是另一个调用虚拟的内核这些对象的方法：

$ cat t1086.cu
#include <stdio.h>
#define N 4


class Polygon {
  protected:
    int width, height;
  public:
  __host__ __device__  void set_values (int a, int b)
      { width=a; height=b; }
  __host__ __device__  virtual int area ()
      { return 0; }
};

class Rectangle: public Polygon {
  public:
  __host__ __device__  int area ()
      { return width * height; }
};

class Triangle: public Polygon {
  public:
  __host__ __device__   int area ()
      { return (width * height / 2); }
};

__global__ void setup_f(Polygon ** d_polys) {
  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  if (idx < N) {
    if (idx%2)
      d_polys[idx] = new Rectangle();
    else
      d_polys[idx] = new Triangle();
    d_polys[idx]->set_values(5,12);
}};

__global__ void area_f(Polygon ** d_polys) {
  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  if (idx < N){
    printf("area of object %d = %d\n", idx, d_polys[idx]->area());
}};


int main () {

  Polygon **devPolys;
  cudaMalloc(&devPolys,N*sizeof(Polygon *));
  setup_f<<<1,N>>>(devPolys);
  area_f<<<1,N>>>(devPolys);
  cudaDeviceSynchronize();
}
$ nvcc -o t1086 t1086.cu
$ cuda-memcheck ./t1086
========= CUDA-MEMCHECK
area of object 0 = 30
area of object 1 = 60
area of object 2 = 30
area of object 3 = 60
========= ERROR SUMMARY: 0 errors
$

Answer 2

罗伯特的建议似乎奏效了：

__global__ void createAProduction(DeviceProduction** production) {
    int idx = threadIdx.x+blockDim.x*blockIdx.x;
    if(idx == 0) {
        production[0] = new AProduction();
    }   
}

这样称呼：

DeviceProduction ** devAProduction = NULL;
cudaMalloc(&devAProduction, sizeof(AProduction *));
createAProduction<<<1, 1>>>(devAProduction);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );

但是，如果我想为 deviceProductions 数组保留单指针结构，是否可以这样做。像这样？

deviceProductions["A"] = (DeviceProduction *) malloc(sizeof(AProduction *));
gpuErrchk(cudaMemcpy(deviceProductions["A"], devAProduction, sizeof(AProduction *), cudaMemcpyDeviceToHost));

我的目的是将指针（地址）从设备内存复制到主机内存。我做得对吗？