Dynamic persistent memory allocation in CUDA

Question

I'm trying to use read-only memory, but it's hard for me to figure out how to nest arrays. I have a data array that takes into account internal data, but they are different for each record. Therefore, based on the following simplified code, I have two problems. At first I don’t know how to distribute the data that the members of my data structure point to. Secondly, since I cannot use cudaGetSymbolAddress for read-only memory, I'm not sure if I can just pass a global pointer (which you cannot do with regular __device__ memory).

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);
}

Thanks for any help. :-)

Answer 1

, 64 , CudaMalloc. , ,

__device__ __constant__ data mydata[100];

, . , , . , segfault ( devicemu).

Answer 2

:

http://forums.nvidia.com/index.php?showtopic=30269&hl=embedded

Answer 3

, .

( 64 ) .

" ", .

Answer 4

"" ? . ,

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

You can simply store this data in an array as follows:

[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)]
]