CUDA /推力双指针问题（指针向量）

Question

嘿，我正在使用CUDA和Thrust库。当我尝试访问CUDA内核上的双指针时遇到问题，该指针使用来自主机的Object *（指针向量）类型的thrust :: device_vector加载。当使用'nvcc -o thrust 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包含主机指针。推力不能在GPU和主机CPU地址空间之间进行“深度复制”或指针转换。因此，当您将spheres_h复制到GPU内存时，您将使用GPU阵列的主机指针进行清理。 GPU上的主机指针的间接指向是非法的 - 它们是指向错误的内存地址空间的指针，因此你在GPU内部获得了与GPU内核相当的segfault。

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

另一种方法将涉及创建一个持有设备指针的主机阵列，然后将该阵列复制到设备内存。你可以在this answer中看到一个例子。请注意，声明包含thrust::device_vector的thrust::device_vector将可能不起作用，因此您可能需要使用基础CUDA API调用来完成此设备指针构造数组。

你也应该注意到我没有提到反向复制操作，这同样很难做到。

底线是推力（和C + + STL容器为此）真的不打算拿指针。它们旨在保存值，并通过使用用户不应该看到的迭代器和底层算法来抽取指针间接和直接访问内存。此外，“深层复制”问题是NVIDIA论坛上的明智人士针对GPU代码中的多级指针提出抗辩的主要原因。它使代码非常复杂，并且在GPU上执行速度也较慢。