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

Question

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

13

使用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++中一样。

我错过了什么？

- JorenHeit

你的目标是拥有一个带有全局函数 _ global _ 的类吗？ - 4pie0

1

那么这个问题的真正目的是为什么CUDA对象模型是这样的，进行一场辩论？这不是一个很好的[SO]问题。投票关闭，因为主要基于观点。 - talonmies

3

我想向Nvidia提出一个功能请求，因为我觉得这很奇怪，竟然不支持这个功能。在我提出请求之前，我想确保没有明显的原因导致它现在的状态。显然没有，但是似乎这是一个否决我的理由。这与意见或辩论无关。 - JorenHeit

1

我们这里没有设计CUDA对象模型的人。因此，我们谁也不能说为什么它不受支持，任何答案都只是推测而已。如果我要猜的话，我会说它是因为它破坏了编译模型 - 一个__global__函数同时在主机对象和设备对象中编译。CUDA类和结构体只能在一个内存空间中实例化。这似乎排除了在结构体和类中有__constant__或__global__对象的可能性。 - talonmies

2

@talonmies：我事先不知道我的问题没有明显的答案，所以我无法预见到猜测。如果你问我，回答/评论中的负面评价和负面语气是非常不必要的。至于你的猜测：当你在类/结构体内声明静态函数时，在内存中不会添加任何内容。声明仅仅为函数添加作用域，类似于命名空间。我将联系NVidia看看他们有什么说法。 - JorenHeit

显示剩余2条评论

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Jakub Klinkovský · Accepted Answer

幸运的是，在这个问题被问出约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却没有。