从静态初始化代码启动CUDA内核时出现问题

Question

从静态初始化代码启动CUDA内核时出现问题

c++cudaglobal-variablesstatic-initialization

11

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

"ScalarField.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");
        }
    };

"classA.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;），那么在内核启动时会出现“无效设备函数”错误。为什么会这样呢？编辑更新了代码以提供更完整的示例。 编辑2 根据Raxvan的评论建议，我想说一下，ScalarField构造函数中使用的dimensions变量也在main函数之外的另一个类中定义，但是在其他所有内容之前定义。这可能是解释吗？虽然调试器显示了正确的dimensions值。

- Noel

你能提供更多的代码来帮助回答这些问题吗：A类在自己的文件中，但内核在另一个文件中，文件扩展名是什么等等。你应该提供足够的代码让其他人能够复制你的问题。 - deathly809

4

如果您将 a_Object 定义为全局变量，它会在全局数据初始化期间开始执行。这是一种非常糟糕的做法，因为无法知道执行顺序。考虑到这一点，初始化所有 CUDA 相关内容的代码可能会晚于全局数据的运行。 - Raxvan

1

@JackOLantern 我修改了代码，但需要同行审查。 - deathly809

@JackOLantern 我注意到它甚至没有进入主方法。我对CUDA初始化过程并不是很熟悉。它必须先进入主方法吗？我尝试使用cudaChooseDevice来初始化CUDA，但仍然无法工作。 - deathly809

“_invalid device function_”通常意味着运行时无法找到与GPU架构相匹配的二进制代码，请参考thrust::device_vector错误。我对为什么不能通过启动内核函数来初始化全局对象没有解释，这可能与C++实例化全局对象的机制有关。肯定有一种解决方法，就是定义一个默认构造函数，可以参考这里。 - Vitality

显示剩余2条评论

1个回答

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

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

- talonmies · Accepted Answer

简短版：

在main函数外部实例化class A时，问题的根本原因是在调用class A的构造函数之前，需要使用特定的钩子程序初始化CUDA运行库与内核程序。这种情况发生是因为C++执行模型中静态对象的初始化顺序无法保证。全局作用域类在CUDA设置对象初始化之前被实例化。在调用内核代码之前，内核代码从未被加载到上下文中，因此会出现运行时错误。

据我所知，这是CUDA运行API的一个真正的限制，不是用户代码容易解决的问题。在您的简单示例中，您可以将内核调用替换为对cudaMemset或其中一种非符号基础的运行时API memset函数的调用，然后它将能够运行。此问题完全限于用户内核或通过运行时API动态加载的设备符号。因此，一个空的默认构造函数也可以解决您的问题。从设计角度来看，我非常怀疑任何在构造函数中调用内核的模式。添加一个特定于GPU的设置/拆卸方法，该方法不依赖于默认构造函数或析构函数，将是一种更清洁和不容易出错的设计。

详细版：

在任何运行时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]

这里是同样的过程，但这次在main内使用了A实例化（这是在执行延迟设置的对象被初始化之后保证发生的）：

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开发者支持并提交错误报告。