代码之家 › 专栏 › 技术社区 › Utku Ufuk

CUDA错误:“\uu device\uuuu、\uu constant\uuuu和\uu shared\uuuu变量不支持动态初始化”

variable-initialization gpu cuda memory-management c++

Utku Ufuk · 技术社区 · 8 年前

我正在尝试静态初始化只读 std::map GPU内存中的变量如下:

// EXAMPLE 1:
using namespace std;

// first attempt: __device__ extern const
__device__ extern const map<char, const char*> BYTES_TO_WORDS = {
{0xB0, "zero"}, {0xB1, "one"}, {0xB2, "two"}, {0xB3, "three"}};

// second attempt: __const__ static
enum class Color{RED, GREEN, BLUE};
enum class Device{PC, TABLET, PHONE};

__constant__ static map<Color, Device> COLORS_TO_THINGS = {
{Color::RED,Device::PC},{Color::GREEN,Device::TABLET},{Color::BLUE,Device::PHONE}};

但我得到以下错误:

dynamic initialization is not supported for __device__, __constant__ and __shared__ variables

我很困惑,因为当我尝试这样的事情时,我没有得到这个错误:

// EXAMPLE 2:
__device__ extern int PLAIN_ARRAY[] = {1, 2, 3, 4, 5};

我只想能够创建和初始化只读 标准::映射 并从CPU和GPU代码访问它。如果你能告诉我怎么做,我将不胜感激。

编辑: 有人指出,设备代码中不支持标准库。但我得到的错误似乎表明这是一个内存管理问题。

1 回复 | 直到 8 年前

Candy Gumdrop 8 年前

初始化C++对象,如 std::map 涉及在运行时调用构造函数。用于初始化 标准::映射 s是 list initialization 这就叫 std::initializer_list 过载 标准::映射 的构造函数。您的示例 PLAIN_ARRAY 不调用任何构造函数,因为这是 aggregate initialization 只需要初始化一些 int ,并初始化 内景 不需要构造函数调用。

在CUDA中,不可能对存储在GPU上的全局变量使用任何类型的动态初始化,例如 __device__ 和 __constant__ 变量,这意味着对象的初始值必须在编译时已知,而不仅仅是在调用构造函数后的运行时生成。

另一个问题是,即使在可以在设备代码中调用构造函数的上下文中,也无法调用 标准::映射 作为C++标准库的一部分,它没有 __设备__ 构造函数,也没有任何其他 __设备__ 成员函数,因此只能从主机代码中使用。CUDA运行时没有为C++STL类定义任何类型的设备功能。即使你设法 cudaMemcpy() 一 标准::映射 从主机内存到GPU内存,您将无法使用该对象,首先是因为它的所有成员函数 __host__ 功能,无 __设备__ 副本,其次是 标准::映射 将在内部包含指针成员变量,该变量引用动态分配的主机内存,该内存将不是GPU上的有效内存地址。

另一种方法是使用结构的普通数组而不是映射,例如:

__device__
const struct {
    unsigned char byte;
    const char word[10];
} BYTES_TO_WORDS[] = {
    {0xB0, "zero"},
    {0xB1, "one"},
    {0xB2, "two"},
    {0xB3, "three"}
};

但是,与 标准::映射 ,您将必须实现通过其键手动查找值。

我只想能够创建和初始化只读 标准::映射 和 从CPU和GPU代码访问它 。

不幸的是,这并不简单,因为您不能将变量定义为两者都是 __设备__ 和 __主机__ 。访问 __设备__ 来自主机代码的变量,您必须使用 cudaMemcpyFromSymbol() ,这与只访问普通变量相比非常尴尬。因此,您可能最终不得不在主机内存中定义常量,然后将常量从主机内存复制到设备内存:

const byte_word BYTES_TO_WORDS[] = {
    {0xB0, "zero"},
    // ...
};

// uninitialized array
__device__
byte_word DEV_BYTES_TO_WORDS[sizeof BYTES_TO_WORDS / sizeof(byte_word)];

// at startup, use `cudaMemCpyToSymbol()` to populate `DEV_BYTES_TO_WORDS`
// from `BYTES_TO_WORDS`.

另一种方法是使用预处理器定义在两个数组中有效地复制和粘贴相同的初始值设定项,而不是在运行时复制数据。在任何情况下,都需要两个单独的阵列。