CUDA 在常量内存中复制数组答案

【问题标题】：CUDA replicating an array in constant memoryCUDA 在常量内存中复制数组
【发布时间】：2012-08-21 10:32:46
【问题描述】：

我有一个映射表，我知道我可以通过执行以下操作将其复制到 CUDA 常量内存：

#define LENGTH 4
#define THREAD_BLOCKS 64

const int mapTable[LENGTH] = {0, 1, 3, 5};

int main()
{
  //..

__constant__ int dMapTable[LENGTH];
cudaMemcpyToSymbol( dMapTable , mapTable, size_t(LENGTH) * sizeof(int) , 0, cudaMemcpyHostToDevice );

  //..
}

现在我要做的是在 CUDA 常量内存中创建该表的多个副本。我要创建的副本数量等于线程块 THREAD_BLOCKS 的数量。谁能建议我如何做到这一点并有效地做到这一点？

【问题讨论】：

这样做的动机是什么？你认为它会提高性能吗？（我对此表示怀疑。）
实际上我内核的所有线程都在常量内存中读取这个数组。到目前为止，我没有使用任何共享内存，computeprof 的统计数据告诉我，大约有 300 万次 warp 执行序列化。此外，根据我的算法，我知道不同的线程将读取此数组中不同且不可预测的位置。通过复制这个数组，我可以摊销由于银行冲突引起的扭曲序列化所产生的成本
为什么不直接将常量复制到共享内存？
是的，我确实考虑过这一点，但我正在考虑将共享内存用于其他内存优化。实际上我仍然不确定在常量内存中复制数组是否可以减少扭曲发散，因为半扭曲中的线程仍然会根据相同的模式访问位置
对于这个应用程序，如果你不能使用共享，L1 缓存可能工作得很好。常量内存针对广播进行了优化，因此绝对不是您在此处寻找的内容。

标签： memory cuda constants

【解决方案1】：

如果您通过设置常量数据的多个副本看到内核性能有任何改进，我会感到非常惊讶。常量内存被缓存了，所以你只会用重复的值来破坏缓存。

此外，值得注意的是，在计算能力 3.0 之前的所有设备上，恒定内存大小仅为 64KiB。

不过，如果您想检查性能，只需像往常一样设置多个副本，然后为内核计时。

【讨论】：

是的。并且常量缓存是small.