Alex Rivera | Logout

Dynamic Allocation of Constant memory in CUDA

Asked 2008-11-07T04:57:35.143
12

I'm trying to take advantage of the constant memory, but I'm having a hard time figuring out how to nest arrays. What I have is an array of data that has counts for internal data but those are different for each entry. So based around the following simplified code I have two problems. First I don't know how to allocate the data pointed to by the members of my data structure. Second, since I can't use cudaGetSymbolAddress for constant memory I'm not sure if I can just pass the global pointer (which you cannot do with plain __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 offered. :-)

Edit
Report

1 Answer

3

No, you cant do that.

Constant memory (64KB max) can only be hard-coded before compilation.

However you can assign texture memory on the fly which is also cached on the Device.

answered 2011-02-25T09:38:03.993

Your Answer