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

Question

我正在向设备上的运行写入一些 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.

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

我在这里错过了什么？

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

谢谢

我在 Ubuntu 14.04 和 gcc 4.8.4

上使用 CUDA7.5 工具包

Answer 1

这道题的重要特点是CUDA使用C++编译模型，dim3被当成class处理。所以同时：

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.

的每个成员

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

How to define constant arrays of dim3 structs in CUDA

struct

cuda

constants

static-initialization