Asignación dinámica de memoria constante en CUDA

Question

Estoy tratando de aprovechar la memoria constante, pero me está costando trabajo encontrar la forma de anidar las matrices. Lo que tengo es una matriz de datos que tiene recuentos para los datos internos pero que son diferentes para cada entrada. Así que, basado en el siguiente código simplificado, tengo dos problemas. Primero, no sé cómo asignar los datos apuntados por los miembros de mi estructura de datos. En segundo lugar, dado que no puedo usar cudaGetSymbolAddress para la memoria constante, no estoy seguro de si puedo pasar el puntero global (que no se puede hacer con la memoria simple __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); 
} 

Gracias por cualquier ayuda ofrecida. :-)

Answer 1

¿Por qué no solo utiliza la denominada representación de datos "empaquetados"? Este enfoque le permite colocar todos los datos que necesita en una matriz de bytes de una dimensión. Por ejemplo, si necesita almacenar

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

Sólo puede almacenar estos datos en la matriz de esta manera:

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

Answer 2

Creo que la memoria constante es de 64 K y no se puede asignar de forma dinámica utilizando CudaMalloc. Debe declararse constante, por ejemplo,

__device__ __constant__ data mydata[100]; 

De forma similar, tampoco es necesario que la libere. Además, no debe pasar la referencia a él mediante un puntero, simplemente acceda a él como una variable global. Intenté hacer algo similar y me dio segfault (en devicemu).

Answer 3

Estas dos hilos debe ayudarle a:

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

Answer 4

No, usted no puede hacer eso.

La memoria constante (64 KB máximo) solo se puede codificar de forma rígida antes de la compilación.

Sin embargo, puede asignar memoria de textura sobre la marcha que también se almacena en caché en el dispositivo.