CUDA：如何为类的数据成员分配内存答案

【问题标题】：CUDA : How to allocate memory for data member of a classCUDA：如何为类的数据成员分配内存
【发布时间】：2012-12-26 10:28:44
【问题描述】：

假设我有这个课程：

class Particle
{
    double *_w;
};

我想将Particle 的nParticles 对象发送到我的内核。为这些对象分配空间很容易：

Particle *dev_p;
cudaStatus = cudaMalloc((void**)&dev_P, nParticles * sizeof(Particle));
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed!");
    goto Error;
}

还假设 nParticles 为 100。现在我需要为Particle 对象中的每个_w 分配300 double。我怎样才能做到这一点？我试过这段代码：

for( int i = 0; i < nParticles; i++){
    cudaStatus = cudaMalloc((void**)&(dev_P[i]._w), 300 * sizeof(double));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }
}

但是当我访问 dev_p[i]._w[j] 时，使用 Nsight 进行调试会停止。

【问题讨论】：

指针dev_P存储了一个设备mem的地址。当您在dev_P[i] 之类的主机上取消引用它时，您正在尝试访问其 addr 值与之前分配的设备 mem 相同的主机 mem。

标签： memory-management cuda

【解决方案1】：

也许您应该包含一个完整的简单示例。（如果我在上面编译您的代码并自行运行它，在 linux 上，我会在第二个 cudaMalloc 操作时遇到 seg 错误）。我看到的一个问题是，由于您在第一步中已在设备内存中分配了粒子对象，因此当您分配 _w 指针时，您将指针传递给已经在设备内存中的 cudaMalloc。您应该将基于主机的指针传递给 cudaMalloc，然后它将分配给设备（全局）内存中的分配区域。

我认为符合我在 yoru 示例中看到的一种可能的解决方案是这样的：

#include <stdio.h>

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

class Particle
{
    public:
    double *_w;
};

__global__ void test(Particle *p){

  int idx=threadIdx.x + blockDim.x*blockIdx.x;

  if (idx == 2){
    printf("dev_p[2]._w[2] = %f\n", p[idx]._w[2]);
    }
}


int main() {
  int nParticles=100;
  Particle *dev_p;
  double *w[nParticles];
  cudaMalloc((void**)&dev_p, nParticles * sizeof(Particle));
  cudaCheckErrors("cudaMalloc1 fail");

  for( int i = 0; i < nParticles; i++){
    cudaMalloc((void**)&(w[i]), 300 * sizeof(double));
    cudaCheckErrors("cudaMalloc2 fail");
    cudaMemcpy(&(dev_p[i]._w), &(w[i]), sizeof(double *), cudaMemcpyHostToDevice);
    cudaCheckErrors("cudaMemcpy1 fail");
    }
  double testval = 32.7;
  cudaMemcpy(w[2]+2, &testval, sizeof(double), cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy2 fail");
  test<<<1, 32>>>(dev_p);
  cudaDeviceSynchronize();
  cudaCheckErrors("kernel fail");
  printf("Done!\n");

}

在这里，我们在主机上创建一组单独的指针以用于 cudaMalloc 目的，然后将这些分配的指针复制到设备以用作设备指针（这对于 UVA 是合法的）。

另一种方法是allocate the _w pointers on the device side。这也可能符合您的目的。

以上所有我都假设 cc 2.0 或更高版本。

使用类似于here 描述的方法，可以将循环中完成的设备端分配折叠为单个分配：

cudaMalloc(&(w[0]), nParticles*300*sizeof(double));
cudaCheckErrors("cudaMalloc2 fail");
cudaMemcpy(&(dev_p[0]._w), &(w[0]), sizeof(double *), cudaMemcpyHostToDevice);
cudaCheckErrors("cudaMemcpy1 fail");
for( int i = 1; i < nParticles; i++){
  w[i] = w[i-1] + 300;
  cudaMemcpy(&(dev_p[i]._w), &(w[i]), sizeof(double *), cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy1 fail");
  }

cudaMemcpy 操作仍需单独完成。

【讨论】：

感谢您提供完整的示例 :)

【解决方案2】：

有两种方法。第一个 - 您在填充 host 粒子对象数组的主机上分配内存。完成后，通过cudaMemcpy 将主机阵列复制到设备。

第二种方式 - 在 Fermi 及更高版本上，您可以在内核中调用 malloc，从内核中填充 dev_P 数组。

【讨论】：