c++ - CUDA 错误 : "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}};

但我收到以下错误:

__device__、__constant__ 和 __shared__ 变量不支持动态初始化

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

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

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

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

Answer 1

初始化 C++ 对象(如 std::map)涉及在运行时调用构造函数。您用于初始化 std::map 的 C++11 语法是 list initialization 的一种形式它调用 std::map 的构造函数的 std::initializer_list 重载。您使用 PLAIN_ARRAY 的示例不会调用任何构造函数，因为这是 aggregate initialization 的一种形式这只涉及按值初始化一些 int，而初始化 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 不同的是，您必须通过键值手动查找值。

---

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`.

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