Allocation dynamique de la mémoire constante dans CUDA

Question

J'essaie de tirer parti de la mémoire constante, mais j'ai du mal à comprendre comment imbriquer des tableaux. Ce que j'ai est un tableau de données qui a des comptes pour les données internes, mais ceux-ci sont différents pour chaque entrée. Donc, basé sur le code simplifié suivant, j'ai deux problèmes. D'abord je ne sais pas comment allouer les données pointées par les membres de ma structure de données. Deuxièmement, puisque je ne peux pas utiliser cudaGetSymbolAddress pour la mémoire constante, je ne suis pas sûr de pouvoir passer le pointeur global (ce que vous ne pouvez pas faire avec la simple mémoire __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); 
} 

Merci pour toute aide offerte. :-)

Answer 1

Pourquoi n'utilisez-vous pas simplement la représentation de données dite "emballée"? Cette approche vous permet de placer toutes les données dont vous avez besoin dans un tableau d'octets à une dimension. Par exemple, si vous devez stocker

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

Vous pouvez simplement stocker ces données dans le tableau de cette façon:

[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

Je pense que la mémoire est constante 64K et vous ne pouvez pas l'affecter dynamiquement à l'aide CudaMalloc. Il doit être déclarée constante, disons,

__device__ __constant__ data mydata[100]; 

De même, vous pouvez aussi ne pas besoin de le libérer. De plus, vous ne devriez pas lui transmettre la référence via un pointeur, il suffit d'y accéder en tant que variable globale. J'ai essayé de faire une chose semblable et ça m'a donné un segfault (in devicemu).

Answer 3

Ces deux fils devraient vous aider:

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

Answer 4

Non, vous ne pouvez pas faire cela.

La mémoire constante (64 Ko maximum) ne peut être codée en dur que lors de la compilation.

Toutefois, vous pouvez affecter une mémoire de texture à la volée qui est également mise en cache sur le périphérique.