在类中声明 cuda __constant__ 内存

Question

我是 CUDA 编程的新手。目前，我正在尝试构建一个使用 CUDA 并行处理数据的 OO 框架。我目前正在使用 CUDA 8.0。

有一些关键参数__constant__ int foo[3]需要所有线程频繁访问，我认为将它们放入常量内存可以帮助改进内存读取过程。我想在类头中声明一个，然后使用cudaMemcpyToSymbol.

从其他示例（不在 OO 上下文中）我了解到可以将常量内存符号定义为全局变量并使用cudaMemcpyToSymbol. 但是，如果我在类的标题中定义常量内存符号，编译器会抱怨：

error : attribute "__constant__" does not apply here

我也不能__constant__在类定义中指定任何地方。貌似在面向对象框架的上下文中无法定义常量内存？如果有任何解决方法，请提供帮助。非常感谢！

---

myClass.cuh

class myClass
{
private:
    __constant__ int foo[3];
    void initialize();
    void compute(); // required repeated and heavy access of foo from every thread
}

myClass.cu

void myClass::initialize()
{
    int bar[3] ={1,2,3};
    //attempt to copy the data in bar to foo
    cudaMemcpyToSymbol(foo, bar, 3*sizeof(int), 0, cudaMemcpyHostToDevice);
}

Answer 1

CUDA 对象模型不允许在结构或类中使用内存说明符 ( __constant__, __shared__, __global__)。这是一个深思熟虑且明智的设计选择——这意味着任何兼容对象都可以在任何内存空间内实例化，并且任何兼容对象都可以在内存空间之间复制。

所以你不能做你所描述的。如果你想要一个小数组存储在常量内存中，那么你必须在你需要访问它的编译单元静态编译它，或者你需要使用单独的编译和链接，在链接中声明一次数组。如果您想要或需要更多的范围控制，则使用命名空间。能做的就这么多。

另一方面，如果您的最终目的是将此结构按值传递给 CUDA 内核，那么您无需担心常量内存。这是因为所有内核参数都是在当前支持的体系结构上的专用常量内存中传递的，您无需执行任何操作即可获得所需的效果。