12

使用 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++ 中习惯的那样。

我错过了什么？

4

1 回答 1

2

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

于 2017-05-14T13:45:36.110 回答