c - 将 CUDA C __constant__ 变量复制到本地内存是否有性能优势

Question

例如，如果您有一个简单的常量变量__device__ __constant__ int MY_CONSTANT;，并且它被同一个内核线程多次访问：

__global__ void move(int* dataA, int* dataB, int* dataC){
    ...
    dataB[threadID] = dataA[threadID] * MY_CONSTANT;
    dataC[threadID] = dataA[[threadID] * dataB[threadID] % MY_CONSTANT;
    ...
}

我可以看到将dataA[threadID]和的值存储dataA[threadID] * MY_CONSTANT在局部变量/寄存器中以避免不必要的全局读取是有益的。忽略这一点，将 的值MY_CONSTANT放在局部变量中以避免它被读取两次是否有益，或者这将由编译器处理，因为它与其他全局数据不同，它不能更改。

Answer 1

当然，唯一知道的方法是对每种方式进行编码并执行以下操作：

1. 生成并检查 PTX 代码以确定编译器为每个实现发出的内容。
2. 使用具有代表性的输入来检测和分析每个实施和测试性能。

您要问的是那种保证“过早优化”这一短语的微优化。我的建议是，除非您发现性能不足以满足您的需求，否则您应该忘记此类优化。

也就是说，常量内存存储在全局内存空间中，当被多处理器访问时，它会被添加到缓存中。对该内存区域的后续访问具有低得多的延迟。当一个扭曲中的所有线程访问相同的常量内存字时，就不会发生冲突，而且一切都会很快发生（请注意，硬件差异很大，我所说的可能在技术上对您的设备并不准确，但是要点是：遵循内存类型的访问模式，您无需担心您在此处询问的内容）。