CUDA:为什么不能定义静态全局成员函数？

Question

当使用nvcc(CUDA 5.0)编译下面的代码时,出现错误"内存限定符的非法组合",因为显然不可能在类中具有全局内核.

class A
{
public:
    __global__ static void kernel();
};

__global__ void A::kernel()
{}

在处理非静态成员时我可以理解这个限制,但是为什么在内核声明为静态时仍然会出现错误？调用此类成员与在函数体中声明函数时调用该函数没有什么不同(A在本例中).

A::kernel <<< 1, 1 >>> ();

有没有理由我错过了为什么还没有实现呢？

编辑:根据答案和评论中的回答,我对我的问题不够清楚.我的问题不是出现错误的原因.显然,这是因为它尚未实施.我的问题是为什么它没有得到实施.到目前为止,我还没有想到让这个功能得以实现的原因.我意识到我可能已经忘记了一个会使问题复杂化的特殊情况,因此也就是问题.

我认为这是一个合理的功能的原因是:

• 静态函数没有this指针所以即使在主机上的对象上调用内核,访问其数据也没有冲突,因为这些数据首先是不可访问的(来自哪个对象的数据？ ).
• 您可能会争辩说,如果类具有与之关联的静态数据,并且存在于主机上,那么原则上应该可以从静态内核访问它.但是,也不支持静态数据,因此不会发生冲突.
• 在主机上的对象上调用静态内核(A a; a.staticKernel<<<...,...>>>();)完全等同于在没有对象的情况下调用它(A::staticKernel<<<...,...>>>();),正如我们在常规C++中习惯的那样.

我错过了什么？

Answer 1

幸运的是，在这个问题提出大约 4 年后，clang 4.0 可以编译 CUDA 语言。考虑这个例子：

class A
{
public:
    __global__ static void kernel();
};

__device__ void A::kernel()
{}

int main()
{
    A::kernel <<< 1, 1 >>> ();
};

当我尝试使用 clang 4.0 编译它时，出现以下错误：

test.cu:7:1: error: kernel function 'kernel' must be a free function or static member function
__global__ void A::kernel()
^
/usr/local/cuda/include/host_defines.h:191:9: note: expanded from macro '__global__'
        __location__(global)
        ^
/usr/local/cuda/include/host_defines.h:88:9: note: expanded from macro '__location__'
        __annotate__(a)
        ^
/usr/local/cuda/include/host_defines.h:86:9: note: expanded from macro '__annotate__'
        __attribute__((a))
        ^
test.cu:7:20: error: __host__ function 'kernel' cannot overload __global__ function 'kernel'
__global__ void A::kernel()
                   ^
test.cu:4:28: note: previous declaration is here
    __global__ static void kernel();
                           ^
2 errors generated.

为了解决这些错误，我将内核定义内联到类声明中：

class A
{
public:
    __global__ static void kernel()
    {
        // implementation would go here
    }
};

然后clang 4.0编译成功，可以执行，没有任何错误。所以这显然不是CUDA语言的限制，而是它事实上的标准编译器。顺便说一句，nvcc 有许多类似的不合理的限制，而 clang 没有。