7

我自己无法弄清楚，确保内核中使用的内存保持不变的最佳方法是什么。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);

这两种方式是等价的还是第一种不保证使用常量内存而不是全局内存？大小动态变化，所以第二种方式在我的情况下并不方便。

4

2 回答 2

14

第二种方法是确保数组编译到 CUDA 常量内存并通过常量内存缓存正确访问的唯一方法。但是您应该问自己，如何在线程块中访问该数组的内容。如果每个线程都会统一访问数组，那么使用常量内存会有性能优势，因为常量内存缓存有广播机制（它也节省了全局内存带宽，因为常量内存存储在片外DRAM和缓存中减少 DRAM 事务计数）。但是如果访问是随机的，那么对本地内存的访问可能会序列化，这将对性能产生负面影响。

可能适合__constant__记忆的典型事物是模型系数、权重和其他需要在运行时设置的常数值。例如，在 Fermi GPU 上，内核参数列表存储在常量内存中。但是，如果内容访问不统一，并且成员的类型或大小在调用之间不是恒定的，那么普通的全局内存更可取。

还要记住，每个 GPU 上下文的常量内存限制为 64kb，因此在常量内存中存储大量数据是不切实际的。如果您需要大量带有缓存的只读存储，则可能值得尝试将数据绑定到纹理并查看性能如何。在 pre-Fermi 卡上，它通常会产生方便的性能提升，在 Fermi 上，与全局内存相比，由于该架构中改进了缓存布局，结果的可预测性可能较低。

于 2012-01-30T06:27:15.910 回答
0

第一种方法将保证函数内部的内存是恒定的Foo。两者不等价，第二个保证它在初始化后保持不变。如果你需要动态而不是你需要使用类似于第一种方式的东西。

于 2012-01-30T04:41:05.560 回答