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

Question

使用 nvcc (CUDA 5.0) 编译下面的代码时，会出现错误“内存限定符的非法组合”，因为在一个类中显然不可能有全局内核。

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

__global__ void A::kernel()
{}

我在处理非静态成员的时候可以理解这个限制，但是为什么内核声明为静态时仍然会出现错误？此类成员的调用与在命名空间中声明函数时调用函数没有什么不同（在这种情况下为A）。

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

我想知道为什么这还没有实施吗？

编辑：根据答案和 cmets 中的回答，我对我的问题还不够清楚。我的问题不是为什么会出现错误。显然，这是因为它尚未实施。我的问题是为什么它没有被实施。到目前为止，我还没有想到阻止此功能实施的原因。我意识到我可能忘记了一个会使事情复杂化的特殊情况，因此提出了这个问题。

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

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

我错过了什么？

Answer 1

幸运的是，在提出这个问题大约 4 年后，clang 4.0 can compile the CUDA language。考虑这个例子：

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 没有。