CUDA代码中的常量内存使用情况

我自己弄不明白，确保内核中使用的内存不变的最佳方法是什么？有一个类似的问题在http://stackoverflow...r-pleasant-way。我正在使用GTX580并仅编译2.0功能。我的内核看起来像CUDA代码中的常量内存使用情况

__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);

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

来源

2012-01-30 markhor

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

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

另外请记住，每个GPU上下文有64kb的常量内存限制，因此在常量内存中存储大量数据是不现实的。如果你需要大量带缓存的只读存储器，可能值得尝试将数据绑定到纹理，并查看性能如何。在费米卡之前，它通常会带来一个方便的性能增益，费米的结果与全局存储器相比，可预测性较差，因为该架构中缓存布局的改进。

来源

2012-01-30 06:27:15 talonmies

第一种方法将保证函数Foo内的内存不变。这两个不相等，第二个保证它在初始化之后在那里存在。如果你需要动态的比你想使用类似于第一种方式的东西。

来源

2012-01-30 04:41:05 rudolph9

CUDA代码中的常量内存使用情况

回答

相关问题