传递指向设备__constant__存储器到内核的指针，而不是使用直接

Question

我使用CUDA 5.0和在Ubuntu 12.10用gcc 4.6 GTX 670

，我已经写入一类称为网格：

https://github.com/benadler/octocopter/blob/master/basestation/grid.cu

https://github.com/benadler/octocopter/blob/master/basestation/grid.cuh

网格级在其它两个类被使用：

• pointcloudcuda
• 粒子系统

现在我想使用Grid的（非静态）方法都pointcloudcuda和粒子系统的内核，即使他们将使用不同的网格（具有不同的价值观不同的网格对象）。因此，对于所有在那里我使用网格类的，我有两个选择：

1）我根本就

Grid hostGrid(...); 
cudaMalloc(gridOnDeviceGlobal, sizeof(Grid)) 
cudaMemcpy(gridOnDeviceGlobal, &hostGrid, sizeof(Grid), cudaMemcpyHostToDevice) 
cloudKernel<<< numBlocks, numThreads >>>(someDate, gridOnDeviceGlobal); 

这很简单，但内核必须从全局存储器读取网格值。这可能会很慢。

2）由于网格值很少改变，我有两个功能

把

__constant__ Grid myGridForPointCloudCuda 

到pointcloudcuda.cu，现在一起

void copyParametersToGpu(Grid *hostGrid) 
{ 
    cudaMemcpyToSymbolAsync(myGridForPointCloudCuda, hostGrid, sizeof(Grid)) 
} 

void getDevicePointerOfGridForPointCloudCuda(Grid** ptr) 
{ 
    cudaGetSymbolAddress((void**)ptr, myGridForPointCloudCuda); 
} 

，在pointcloudcuda.cpp，我可以

Grid hostGrid(...); 
copyParametersToGpu(&hostGrid); 
Grid* gridOnDeviceConstant; 
getDevicePointerOfGridForPointCloudCuda(&gridOnDeviceConstant); 
cloudKernel<<< numBlocks, numThreads >>>(someDate, gridOnDeviceConstant); 

在我看来，2）的优势将是更快访问内核中的常量内存。然而，在其他地方，我读到这不起作用，因为编译器编译CUDA内核在编译时并不知道传递的网格指针是指向全局还是常量内存，因此必须使用较慢内存取指令。

威尔2）是在GeForce GTX的670比1快）？

有没有更好的办法做到我想要什么？我只需要将不同的Grid实例传递给内核。在开始使用多个Grid实例之前，常量变量是一个舒适且快速的选择。

谢谢！

Answer 1

如果有多个网格实例不是简单地分配在恒定存储器中的栅格阵列，网格实例复制到阵列和调用当内核通过一个指数到网格阵列，而不是一个指针指向一个特定网格实例。在内核中使用索引来访问特定的网格实例。