CUDA/Thrust 双指针问题（指针向量）

Question

大家好，我正在使用 CUDA 和 Thrust 库。当我尝试访问 CUDA 内核上的双指针时遇到问题，该内核加载了来自主机的 Object*（指针向量）类型的推力::device_vector。使用“nvcc -o 推力 main.cpp cukernel.cu”编译时，我收到警告“警告：无法判断指针指向的内容，假设为全局内存空间”，并在尝试运行程序时出现启动错误。

我已阅读 Nvidia 论坛，解决方案似乎是“不要在 CUDA 内核中使用双指针”。在发送到内核之前，我不希望将双指针折叠成一维指针......有没有人找到解决这个问题的方法？所需代码如下，在此先感谢！

--------------------------
        main.cpp
--------------------------

Sphere * parseSphere(int i)
{
  Sphere * s = new Sphere();
  s->a = 1+i;
  s->b = 2+i;
  s->c = 3+i;
  return s;
}

int main( int argc, char** argv ) {

  int i;
  thrust::host_vector<Sphere *> spheres_h;
  thrust::host_vector<Sphere> spheres_resh(NUM_OBJECTS);

  //initialize spheres_h
  for(i=0;i<NUM_OBJECTS;i++){
    Sphere * sphere = parseSphere(i);
    spheres_h.push_back(sphere);
  }

  //initialize spheres_resh
  for(i=0;i<NUM_OBJECTS;i++){
    spheres_resh[i].a = 1;
    spheres_resh[i].b = 1;
    spheres_resh[i].c = 1;
  }

  thrust::device_vector<Sphere *> spheres_dv = spheres_h;
  thrust::device_vector<Sphere> spheres_resv = spheres_resh;
  Sphere ** spheres_d = thrust::raw_pointer_cast(&spheres_dv[0]);
  Sphere * spheres_res = thrust::raw_pointer_cast(&spheres_resv[0]);

  kernelBegin(spheres_d,spheres_res,NUM_OBJECTS);

  thrust::copy(spheres_dv.begin(),spheres_dv.end(),spheres_h.begin());
  thrust::copy(spheres_resv.begin(),spheres_resv.end(),spheres_resh.begin());

  bool result = true;

  for(i=0;i<NUM_OBJECTS;i++){
    result &= (spheres_resh[i].a == i+1);
    result &= (spheres_resh[i].b == i+2);
    result &= (spheres_resh[i].c == i+3);
  }

  if(result)
  {
    cout << "Data GOOD!" << endl;
  }else{
    cout << "Data BAD!" << endl;
  }

  return 0;
}


--------------------------
        cukernel.cu
--------------------------
__global__ void deviceBegin(Sphere ** spheres_d, Sphere * spheres_res, float    
num_objects)
{
  int index = threadIdx.x + blockIdx.x*blockDim.x;

  spheres_res[index].a = (*(spheres_d+index))->a; //causes warning/launch error
  spheres_res[index].b = (*(spheres_d+index))->b; 
  spheres_res[index].c = (*(spheres_d+index))->c; 
}

void kernelBegin(Sphere ** spheres_d, Sphere * spheres_res, float num_objects)
{

 int threads = 512;//per block
 int grids = ((num_objects)/threads)+1;//blocks per grid

 deviceBegin<<<grids,threads>>>(spheres_d, spheres_res, num_objects);
}

Answer 1

这里的基本问题是设备向量spheres_dv 包含主机指针。 Thrust 不能在 GPU 和主机 CPU 地址空间之间进行“深度复制”或指针转换。因此，当您将spheres_h 复制到 GPU 内存时，您将得到一个 GPU 主机指针数组。 GPU 上主机指针的间接寻址是非法的 - 它们是错误内存地址空间中的指针，因此您在内核中获得了相当于 GPU 的段错误。

解决方案将涉及用在 GPU 上执行内存分配的东西替换您的 parseSphere 函数，而不是使用目前在主机内存中分配每个新结构的 parseSphere。如果您有一个 Fermi GPU（看起来您没有）并且使用的是 CUDA 3.2 或 4.0，那么一种方法是将parseSphere 转换为内核。设备代码支持 C++ new 运算符，因此结构创建将在设备内存中进行。您需要修改Sphere 的定义，以便将构造函数定义为__device__ 函数才能使这种方法起作用。

另一种方法是创建一个包含设备指针的主机数组，然后将该数组复制到设备内存中。你可以在this answer 中看到一个例子。请注意，声明包含 thrust::device_vector 的 thrust::device_vector 可能无法正常工作，因此您可能需要使用底层 CUDA API 调用来构建这个设备指针数组。

您还应该注意，我没有提到反向复制操作，这同样难以做到。

最重要的是，thrust（以及 C++ STL 容器）确实不是用来保存指针的。它们旨在保存值，并通过使用用户不应该看到的迭代器和底层算法抽象出指针间接和直接内存访问。此外，“深拷贝”问题是 NVIDIA 论坛上的智者反对 GPU 代码中的多级指针的主要原因。它使代码变得非常复杂，并且在 GPU 上执行速度也较慢。