cuda - 如何在主机和设备代码中最好地利用相同的常量？

Question

假设我有一些在主机端代码中使用的全局常量数据：

const float my_array[20] = { 45.146, 54.633, 74.669, 12.734, 74.240, 100.524 };

（注意：我把它们保留为 C-ish，constexpr这里没有。）

我现在也想在设备端代码中使用这些。我不能简单地开始使用它们：它们不能从设备直接访问，并且尝试使用它们会给出：

error: identifier "my_array" is undefined in device code 

使这些常量在主机和设备上都可用的惯用方式是什么？

Answer 1

Mark Harris在 2012 年的回答中提出了这种方法：

#define MY_ARRAY_VALUES 45.146, 54.633, 74.669, 12.734, 74.240, 100.524
__constant__ float device_side_my_array[2] = { MY_ARRAY_VALUES };
const        float host_side_my_array[2]   = { MY_ARRAY_VALUES };
#undef MY_ARRAY_VALUES

__device__ __host__ float my_array(size_t i) {
#ifdef __CUDA_ARCH__
    return device_side_my_array[i];
#else
    return host_side_my_array[i];
#endif
}

但这有一些缺点：

• 实际上没有使用相同的常量，只是具有相同值的常量。
• 数据重复。
• 占用恒定内存，这是相当有限的资源。
• 似乎有点冗长（尽管其他选项可能更是如此）。

我想知道这是否是大多数人在实践中使用的。

笔记：

• 在 C++ 中，可能使用相同的名称，但在命名空间内的不同子命名detail::空间中。
• 不使用cudaMemcpyToSymbol()。