CUDA代码中的常量内存使用

Question

我自己想不出来，如何最好地确保我的内核中使用的内存是恒定的。在http://stackoverflow...r-pleasant-way上有一个类似的问题。

我正在使用GTX580，并且仅编译2.0能力。我的内核看起来像：

__global__ Foo(const int *src, float *result) {...}

我在主机上执行以下代码：

cudaMalloc(src, size);
cudaMemcpy(src, hostSrc, size, cudaMemcpyHostToDevice);
Foo<<<...>>>(src, result);

另一种方法是添加。

__constant__ src[size];

将代码转换为 .cu 文件，从内核中删除 src 指针并执行。

cudaMemcpyToSymbol("src", hostSrc, size, 0, cudaMemcpyHostToDevice);
Foo<<<...>>>(result);

这两种方法等效吗？第一种方法是否不能保证使用常量内存而不是全局内存？由于size是动态变化的，所以第二种方法在我的情况下不方便。

- markhor

Answer 1

第二种方法是确保数组编译到CUDA常量内存并通过常量内存缓存正确访问的唯一方法。但是，你应该考虑在一个线程块内如何访问该数组的内容。如果每个线程都以相同方式访问数组，则使用常量内存将有性能优势，因为常量内存缓存中存在广播机制（它还节省了全局内存带宽，因为常量内存存储在片外DRAM中，并且缓存减少了DRAM事务计数）。但如果访问时是随机的，则可能会出现对局部内存的串行访问，这将对性能产生负面影响。

适合__constant__内存的典型内容包括模型系数、权重和其他需要在运行时设置的常量值。例如，在Fermi GPU上，内核参数列表存储在常量内存中。但如果成员的类型或大小不是从调用到调用都恒定的，或者内容的访问是非均匀的，则通常首选普通的全局内存。

此外，请记住，每个GPU上下文的常量内存限制为64kb，因此在常量内存中存储大量数据是不实际的。如果需要大量只读存储并具有缓存功能，可以尝试将数据绑定到纹理并查看其性能如何。在Fermi卡之前，它通常会产生方便的性能增益，在Fermi上与全局内存相比结果可能不太可预测，因为该架构中缓存布局有所改进。

- talonmies

Answer 2

第一种方法将确保函数Foo内存是恒定的。两者并不相等，第二种方法保证其在初始化后恒定。如果您需要动态，则需要使用类似于第一种方式的东西。

- rudolph9