CUDA中常量内存的动态分配
Posted
技术标签:
【中文标题】CUDA中常量内存的动态分配【英文标题】:Dynamic Allocation of Constant memory in CUDA 【发布时间】:2010-09-21 05:45:00 【问题描述】:我正在尝试利用常量内存,但我很难弄清楚如何嵌套数组。我所拥有的是一组数据,其中包含内部数据的数量,但每个条目的数量都不同。因此,基于以下简化代码,我有两个问题。首先我不知道如何分配我的数据结构成员指向的数据。其次,由于我不能将 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);
感谢您提供的任何帮助。 :-)
【问题讨论】:
【参考方案1】:我认为常量内存是 64K,您不能使用 cudaMalloc
动态分配它。它必须被声明为常量,比如说,
__constant__ data mydata[100];
同样你也不需要释放它。此外,您不应该通过指针传递对它的引用,只需将其作为全局变量访问即可。我尝试做类似的事情,它给了我段错误(在 devicemu 中)。
【讨论】:
【参考方案2】:不,你不能那样做。
恒定内存(最大 64KB)只能在编译前进行硬编码。
但是,您可以动态分配纹理内存,该内存也缓存在设备上。
【讨论】:
【参考方案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)]
]
【讨论】:
以上是关于CUDA中常量内存的动态分配的主要内容,如果未能解决你的问题,请参考以下文章