Dynamiczne przydzielanie stałej pamięci w CUDA

Question

Próbuję skorzystać ze stałej pamięci, ale ciężko jest mi zastanowić się, jak zagnieździć tablice. To, co mam, to tablica danych, która liczy się dla danych wewnętrznych, ale są one różne dla każdego wpisu. Opierając się na następującym uproszczonym kodzie, mam dwa problemy. Najpierw nie wiem, jak przydzielić dane wskazywane przez członków mojej struktury danych. Po drugie, ponieważ nie mogę używać cudaGetSymbolAddress do stałej pamięci, nie jestem pewien, czy mogę po prostu przekazać globalny wskaźnik (którego nie można zrobić zwykłą pamięcią __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); 
} 

Dzięki za pomoc. :-)

Answer 1

Dlaczego po prostu nie używasz tak zwanej "upakowanej" reprezentacji danych? To podejście umożliwia umieszczenie wszystkich potrzebnych danych w jednowymiarowej tablicy bajtów. Na przykład, jeśli trzeba przechowywać

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

Wystarczy przechowywać te dane w tablicy w ten sposób:

[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

myślę stały 64K pamięci i nie można przeznaczyć go dynamicznie przy użyciu CudaMalloc. Musi być zadeklarowany jako stały, na przykład:

__device__ __constant__ data mydata[100]; 

Podobnie ty również nie musisz go uwolnić. Nie powinieneś też przekazywać referencji do niego za pomocą wskaźnika, wystarczy uzyskać do niego dostęp jako zmienną globalną. Próbowałem zrobić coś podobnego i dało mi to segfault (w devicemu).

Answer 3

Te dwa wątki powinny pomóc:

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

Answer 4

Nie, nie można tego zrobić.

Stałą pamięć (maks. 64 KB) można zakodować tylko przed kompilacją.

Można jednak przypisać pamięć tekstury w locie, która jest również buforowana na urządzeniu.