如何在 CUDA 中定义 dim3 结构的常量数组答案

【问题标题】：How to define constant arrays of dim3 structs in CUDA如何在 CUDA 中定义 dim3 结构的常量数组
【发布时间】：2016-07-10 16:13:19
【问题描述】：

我正在编写一些 CUDA 代码以在设备上运行。该代码将使用两个常量值查找表。其中第一个是 256 个无符号整数的数组，我将其声明为：

__constant__ 
uint16_t edgeTable[256]={
   0x000,
   0x019,
   ... etc.
};

这似乎编译得很好。

第二个是dim3的固定大小数组，我试过这个：

__constant__
dim3 offsets[8] = {
    {0, 0, 0}, {0, 0, 1}, {0, 1, 0},
    ... etc 
};

编译器反对的对象。带有错误消息：

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

也许我误解了动态初始化，但在我看来这是静态初始化，编译器可以计算出所有内容的大小并提供所有值。

我在这里错过了什么？

我怎样才能实现我想要做的事情？

谢谢

我在带有 gcc 4.8.4 的 Ubuntu 14.04 上使用 CUDA7.5 工具包

【问题讨论】：

可能是个愚蠢的问题——你是否传递了启用 C++11 的标志？
在 C++11 之前的 C++ 中不支持参数化构造函数初始化（dim3 是一个类）
也许你的问题的答案是，这实际上是动态初始化的一种形式（参见上面的 cmets），并且这些变量类型不支持动态初始化，正如编译器所指出的那样，无论 c ++11 编译器开关。作为一种可能的解决方法，您可以创建一个int 数组（__constant__ int 的静态初始化），然后在内核代码中将其转换为dim3 数组，可能类似于this。我认为您可能已经自己想出了这样的解决方法，所以，我并不是说这是一个答案。
卫生署！谢谢。我确实通过了 -std=c++11 但是完全没有考虑结构是一个类的事实，因此当然会被动态初始化。我认为@Robert Crovella 你的回答是我唯一能做的..

标签： struct cuda constants static-initialization

【解决方案1】：

这个问题的重要特点是CUDA使用了C++编译模型，dim3被当作一个类来对待。所以虽然：

dim3 foo = {1,1,1};

在 C++11 中是合法的，因为参数化构造函数初始化支持，这个：

__constant__ dim3 foo = {1,1,1};

不是，因为这意味着常量内存对象的动态初始化，而 CUDA 执行模型不允许这样做。

如果常量内存方面对您很重要，并且您想要dim3 的便利，您可以这样做：

#include <cstdio>

__constant__ int offsets[3*8];

__global__ void kernel()
{
    if (threadIdx.x < 8) {
        dim3 val = *reinterpret_cast<dim3*>(&offsets[3*threadIdx.x]);
        printf("%d (%d,%d,%d)\n", threadIdx.x, val.x, val.y, val.z);
    }
}

void setup_offsets()
{
    // This requires C++11 support
    dim3 ovals[8] = { {0,0,0}, 
                      {1,0,0}, {0,1,0}, {0,0,1},
                      {1,1,0}, {1,0,1}, {0,1,1},
                      {1,1,1} };

    cudaMemcpyToSymbol(offsets, &ovals[0], sizeof(ovals));
}

int main(void)
{
    setup_offsets();
    kernel<<<1,8>>>();
    cudaDeviceSynchronize();
    cudaDeviceReset();
    return 0;
}

这有点老套，但可能是您在这种情况下所能期望的最好的。查看该代码的 PTX，编译器已正确发出 ld.const.u32 以获取 dim3 的每个成员。

【讨论】：