CUDA：内存限定符的非法组合答案

【问题标题】：CUDA: illegal combination of memory qualifiersCUDA：内存限定符的非法组合
【发布时间】：2016-12-12 22:22:09
【问题描述】：

我有以下代码：

main.cu:

#include "class.h"
int main () {}

class.h:

class Class {
    __global__
    void Function() {};
};

当我使用命令nvcc -c main.cu -o main.o 编译此代码时，我收到以下错误：

class.h(3): warning: inline qualifier ignored for "global" function
class.h(3): error: illegal combination of memory qualifiers

我对这些错误中的每一个都有疑问。为什么它“忽略”函数的__global__ 限定符，为什么__global__ 内存限定符在这种情况下是非法的？我在documentation 中读到了

E.2.10.2. Function Members
Static member functions cannot be __global__ functions.

但是，据我所知，我的函数不是静态成员。删除__global__ 行可以编译，将__global__ 和void Function(); 行移到main.cu 中也是如此。如果这实际上是不允许的，为什么 CUDA 会强制执行此限制，有什么方法可以在保持结构化代码的同时解决这个问题？

为了澄清，我知道没有其他方法可以制作具有可以创建 GPU 内核的函数的类。在我看来，内核只能从 main.cu 中的全局函数创建。我对 CUDA 编程相当陌生，所以我可能只是错过了一些我可能不清楚的 CUDA 约定。如果是这种情况，请告诉我，以便我能跟上正确的编程习惯。

【问题讨论】：

简短的回答是不，你不能这样做。如果你用谷歌搜索“cuda global class member”，你会发现很多关于这个的处理方法，包括像here和here这样的SO问题，你的问题可以说是这些问题的重复。作为一个简单的建议，您可以将 cuda 内核包装在主机可调用的类成员函数中，以“跟上正确的编程实践”。
您好@Robert，感谢您的评论。我只是想确保我理解你的建议。您是否建议我在我的班级中创建一个__host__function，并且该函数的实现调用一个__global__ 函数？如果是这样，__global__ 函数应该在什么范围内？
是的，第一个问题。不确定我是否理解第二个问题。类成员函数的实现应该在什么范围内是否有任何不清楚的地方？
好的，我现在可能明白了。 __global__ 限定符不独立于全局范围的概念。也就是说，任何__global__ 函数必须是全局函数，因为它是在全局范围内定义的。我的理解是__global__ 严格意味着该函数是在主机和设备上定义的，而全局范围内的函数可以从主机上的任何对象访问。如果此解释正确，请将您的答案作为正式答案发布，以便我接受。

标签： c++ cuda

【解决方案1】：

我的理解是您希望以 OOP 方式使用 CUDA 内核。如果是这种情况，下面的类结构应该可以工作：

// myclass.h
class MyClass {
    public:
        void call_kernel( ... );
};

// myclass.cu
__global__
void my_kernel( ... ) {
    // do some work
}

void MyClass::call_kernel() {
    // prepare data for the kernel, e.g. allocating memory, copying from host to device, etc.

    // run kernel
    my_kernel <<< ... >>>( ... );

    // copy results from device to host, clean up, etc.
}

请注意，如果您有多个包含内核代码的类，它们的源代码文件都应该使用.cu 扩展名，并且您应该启用separate compilation。

【讨论】：