CUDA 错误：“__device__、__constant__ 和 __shared__ 变量不支持动态初始化”答案

【问题标题】：CUDA error: "dynamic initialization is not supported for __device__, __constant__ and __shared__ variables"CUDA 错误：“__device__、__constant__ 和 __shared__ 变量不支持动态初始化”
【发布时间】：2018-09-15 11:21:38
【问题描述】：

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

// 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};

我只想能够创建和初始化一个只读 std::map 并从 CPU 和 GPU 代码访问它。如果您能告诉我如何正确操作，我将不胜感激。

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

【问题讨论】：

std::map 将无法从 GPU 代码中正确使用，即使您可以解决此存储问题。限制标识为here。
谢谢，我不知道。
错误很明显。 PLAIN_ARRAY 和 BYTES_TO_WORDS 的初始化过程不一样，因为一个是静态的，另一个是动态的，并且涉及运行构造函数，这对于静态声明的设备符号是非法的

标签： c++ memory-management cuda gpu variable-initialization

【解决方案1】：

初始化诸如std::map 之类的C++ 对象涉及在运行时调用构造函数。您用来初始化std::maps 的C++11 语法是list initialization 的一种形式，它调用std::map 的构造函数的std::initializer_list 重载。您的 PLAIN_ARRAY 示例不调用任何构造函数，因为这是 aggregate initialization 的一种形式，它只涉及按值初始化一些 ints，而初始化 int 不需要构造函数调用。

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

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

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

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

但是，与std::map 不同，您必须手动通过键查找值。

我只是希望能够创建和初始化只读std::map 并从 CPU 和 GPU 代码访问它。

不幸的是，这不是微不足道的，因为您不能将变量定义为__device__ 和__host__。要从主机代码访问__device__ 变量，您必须使用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`.

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

【讨论】：