CUDA error: "dynamic initialization is not supported for device, constant and shared variables"

Question

我正在尝试静态初始化 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 代码访问它。如果你能告诉我如何正确地做，我将不胜感激。

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

Answer 1

初始化 C++ 对象（例如 std::map）涉及在运行时调用构造函数。您用于初始化 std::maps 的 C++11 语法是 list initialization which calls the std::initializer_list overload of std::map's constructor. Your example with PLAIN_ARRAY does not call any constructors as this is a form of aggregate initialization 的一种形式，它只涉及按值初始化一些 ints，而初始化 int 确实不需要构造函数调用。

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

另一个问题是，即使在您可以在设备代码中调用构造函数的上下文中，您也无法调用 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 不同的是，您必须通过键值手动查找值。

I just want to be able to create and initialize a read-only std::map and access it from both CPU and GPU code.

不幸的是，这不是微不足道的，因为您不能将变量定义为 __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`.

另一种方法是使用预处理器定义来有效地跨两个数组复制和粘贴相同的初始值设定项，而不是在运行时复制数据。无论如何，都需要两个单独的数组。

CUDA error: "dynamic initialization is not supported for device, constant and shared variables"