将指向设备 __constant__ 内存的指针传递给内核，而不是直接使用答案

【问题标题】：Passing a pointer that points to device __constant__ memory to kernels, instead of using directly将指向设备 __constant__ 内存的指针传递给内核，而不是直接使用
【发布时间】：2013-02-15 03:05:12
【问题描述】：

我在带有 gcc 4.6 的 ubuntu 12.10 上使用 CUDA 5.0 和 GTX 670，并且我编写了一个名为 Grid 的类：

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

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

Grid 类被用于另外两个类：

pointcloudcuda
粒子系统

现在我想在 pointcloudcuda 和particlesystem 的内核中使用网格的（非静态）方法，即使它们将使用不同的网格（具有不同值的不同网格对象）。因此，对于我使用 Grid 的所有类，我有两种选择：

1) 我只是这样做

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

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

2) 由于 Grid-values 很少改变，我放了一个

__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 内核的编译器在编译时不知道传递的网格指针是指向全局内存还是常量内存，因此必须使用较慢内存获取指令。

在 Geforce GTX 670 上 2) 会比 1) 快吗？

有没有更好的方法来做我想做的事？我只需要将不同的网格实例传递给内核。在我开始使用多个 Grid 实例之前，常量变量是一种既方便又快捷的选择。

谢谢！

【问题讨论】：

标签： pointers memory cuda constants

【解决方案1】：

如果您有多个 Grid 实例而不是简单地在常量内存中分配一个 Grid Array，请将 Grid 实例复制到该数组中，并在调用内核时将一个索引传递给 Grid Array，而不是指向特定 Grid 实例的指针。在内核内部使用索引来访问特定的 Grid 实例。

【讨论】：