CUDA中常量内存的动态分配

Question

我试图利用常量内存，但是我很难弄清楚如何嵌套数组。我有一个数据数组，其中包含内部数据的计数，但每个条目的计数都不同。 基于以下简化代码，我有两个问题。首先，我不知道如何分配由我的数据结构成员指向的数据。其次，由于无法对常量内存使用cudaGetSymbolAddress，我不确定是否可以仅传递全局指针（这在普通__device__内存中是不允许的）。

struct __align(16)__ data{
int nFiles;
int nNames;
int* files;
int* names;
};

__device__ __constant__ data *mydata;

__host__ void initMemory(...)
{
    cudaMalloc( (void **) &(mydata), sizeof(data)*dynamicsize );
    for(int i=; i lessthan dynamicsize; i++)
    {
        cudaMemcpyToSymbol(mydata, &(nFiles[i]), sizeof(int), sizeof(data)*i, cudaMemcpyHostToDevice);
        //...
        //Problem 1: Allocate & Set mydata[i].files
    }
}

__global__ void myKernel(data *constDataPtr)
{
    //Problem 2: Access constDataPtr[n].files, etc
}

int main()
{
    //...
    myKernel grid, threads (mydata);
}

感谢提供任何帮助。:-)

- Ben Tribelhorn

Answer 1

我认为常量内存是64K，您不能使用cudaMalloc动态分配它。它必须被声明为常量，例如：

__constant__ data mydata[100];

同样的，您也不需要释放它。此外，您不应该通过指针传递对它的引用，只需将其作为全局变量访问即可。我尝试过类似的事情，但在devicemu中导致了segfault。

- Nishu

Answer 2

不可以这样做。

常量内存（最大64KB）只能在编译之前硬编码。

然而，您可以在运行时分配纹理内存，该内存也会被缓存在设备上。

- jwdmsd

Answer 3

为什么不使用所谓的“紧凑”数据表示法呢？这种方法允许您将所有需要的数据放入一维字节数组中。例如，如果您需要存储

struct data
{
    int nFiles;
    int nNames;
    int* files;
    int* names;
}

你可以这样将数据存储在数组中：

[struct data (7*4=28 bytes)
    [int nFiles=3 (4 bytes)]
    [int nNames=2 (4 bytes)]
    [file0 (4 bytes)]
    [file1 (4 bytes)]
    [file2 (4 bytes)]
    [name0 (4 bytes)]
    [name1 (4 bytes)]
]

- izogfif