无法从静态初始化代码启动 CUDA 内核

我有一个在其构造函数中调用内核的类，如下所示：

“标量场.h”

#include <iostream>

    void ERROR_CHECK(cudaError_t err,const char * msg) {
        if(err!=cudaSuccess) {
            std::cout << msg << " : " << cudaGetErrorString(err) << std::endl;
            std::exit(-1);
        }
    }

    class ScalarField {
    public:
        float* array;
        int dimension;

        ScalarField(int dim): dimension(dim) {
            std::cout << "Scalar Field" << std::endl;
            ERROR_CHECK(cudaMalloc(&array, dim*sizeof(float)),"cudaMalloc");
        }
    };

“A.h 级”

#include "ScalarField.h"


static __global__ void KernelSetScalarField(ScalarField v) {
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    if (index < v.dimension) v.array[index] = 0.0f;
}

class A {
public:
    ScalarField v;

    A(): v(ScalarField(3)) {
        std::cout << "Class A" << std::endl;
        KernelSetScalarField<<<1, 32>>>(v);
        ERROR_CHECK(cudaGetLastError(),"Kernel");
    }
};

“main.cu”

#include "classA.h"

A a_object;

int main() {
    std::cout << "Main" << std::endl;
    return 0;
}

如果我在 main 上实例化这个类（A a_object;）我没有收到任何错误。但是，如果我在 main 之外实例化它，就在定义它之后（class A {...} a_object;）当内核启动时，我收到“无效设备功能”错误。为什么会发生这种情况？

EDIT

更新了代码以提供更完整的示例。

EDIT 2

根据 Raxvan 评论中的建议，我想说我有dimensionsScalarField 构造函数中使用的变量也在 main 之外（但在其他所有内容之前）定义（在另一个类中）。可以这样解释吗？调试器显示了正确的值dimensions though.

简短版本：

问题的根本原因是什么时候class A在 main 之外实例化的原因是，使用内核初始化 CUDA 运行时库所需的特定钩子例程在构造函数之前没有运行class A正在被呼叫。发生这种情况是因为无法保证 C++ 执行模型中静态对象实例化和初始化的顺序。在初始化执行 CUDA 设置的全局范围对象之前，将实例化您的全局范围类。您的内核代码在调用之前永远不会加载到上下文中，并且会导致运行时错误。

据我所知，这是 CUDA 运行时 API 的真正限制，并且在用户代码中不容易修复。在您的简单示例中，您可以将内核调用替换为对cudaMemset或非基于符号的运行时 API memset 函数之一，它将起作用。这个问题完全仅限于通过运行时 API 在运行时加载的用户内核或设备符号。因此，一个空的默认构造函数也可以解决您的问题。从设计的角度来看，我对任何在构造函数中调用内核的模式都非常怀疑。恕我直言，为类 GPU 安装/拆卸添加一个不依赖于默认构造函数或析构函数的特定方法将是一个更干净且不易出错的设计。

详细地：

有一个内部生成的例程（__cudaRegisterFatBinary）必须运行该命令以使用 CUDA 驱动程序 API 加载和注册任何运行时 API 程序的 fatbin 有效负载中包含的内核、纹理和静态定义的设备符号，然后才能正确调用内核。这是运行时 API 的“惰性”上下文初始化功能的一部分。您可以通过以下方式自行确认：

这是您发布的修改后的示例的 gdb 跟踪。注意我插入了一个断点__cudaRegisterFatBinary，并且在你的静态之前没有达到A构造函数被调用并且内核启动失败：

talonmies@box:~$ gdb a.out 
GNU gdb (Ubuntu/Linaro 7.4-2012.04-0ubuntu2.1) 7.4-2012.04
Copyright (C) 2012 Free Software Foundation, Inc.
License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html>
This is free software: you are free to change and redistribute it.
There is NO WARRANTY, to the extent permitted by law.  Type "show copying"
and "show warranty" for details.
This GDB was configured as "x86_64-linux-gnu".
For bug reporting instructions, please see:
<http://bugs.launchpad.net/gdb-linaro/>...
Reading symbols from /home/talonmies/a.out...done.
(gdb) break '__cudaRegisterFatBinary' 
Breakpoint 1 at 0x403180
(gdb) run
Starting program: /home/talonmies/a.out 
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
Scalar Field
[New Thread 0x7ffff5a63700 (LWP 10774)]
Class A
Kernel : invalid device function 
[Thread 0x7ffff5a63700 (LWP 10774) exited]
[Inferior 1 (process 10771) exited with code 0377]

这是相同的过程，这次是A内部实例化main（这肯定会在执行延迟设置的对象初始化后发生）：

talonmies@box:~$ cat main.cu
#include "classA.h"


int main() {
    A a_object;
    std::cout << "Main" << std::endl;
    return 0;
}

talonmies@box:~$ nvcc --keep -arch=sm_30 -g main.cu
talonmies@box:~$ gdb a.out 
GNU gdb (Ubuntu/Linaro 7.4-2012.04-0ubuntu2.1) 7.4-2012.04
Copyright (C) 2012 Free Software Foundation, Inc.
License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html>
This is free software: you are free to change and redistribute it.
There is NO WARRANTY, to the extent permitted by law.  Type "show copying"
and "show warranty" for details.
This GDB was configured as "x86_64-linux-gnu".
For bug reporting instructions, please see:
<http://bugs.launchpad.net/gdb-linaro/>...
Reading symbols from /home/talonmies/a.out...done.
(gdb) break '__cudaRegisterFatBinary' 
Breakpoint 1 at 0x403180
(gdb) run
Starting program: /home/talonmies/a.out 
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".

Breakpoint 1, 0x0000000000403180 in __cudaRegisterFatBinary ()
(gdb) cont
Continuing.
Scalar Field
[New Thread 0x7ffff5a63700 (LWP 11084)]
Class A
Main
[Thread 0x7ffff5a63700 (LWP 11084) exited]
[Inferior 1 (process 11081) exited normally]

如果这对您来说确实是一个严重的问题，我建议联系 NVIDIA 开发人员支持并提出错误报告。