CUDA alloca memoria nella funzione __device__

Question

C'è un modo in CUDA di allocare memoria nella funzione __device__? Non sono riuscito a trovare alcun esempio di ciò.

Dal manuale: B.15 Allocazione memoria globale dinamica void * malloc (size_t size); libero da void (void * ptr); allocare e liberare memoria dinamicamente da un heap di dimensioni fisse nella memoria globale. La funzione malloc nel kernel CUDA alloca almeno i byte di dimensione dall'heap della periferica e restituisce un puntatore alla memoria allocata o NULL se la memoria insufficiente esiste per soddisfare la richiesta. Il puntatore restituito è garantito per essere allineato a un limite di 16 byte. La funzione CUDA in-kernel free() rilascia la memoria puntata da ptr, che deve essere stata restituita da una precedente chiamata a malloc(). Se ptr è NULL, la chiamata a free() viene ignorata. Chiamate ripetute a free() con lo stesso ptr ha un comportamento indefinito. La memoria allocata da un determinato thread CUDA tramite malloc() rimane allocata per la durata del contesto CUDA o finché non viene rilasciata esplicitamente da una chiamata a free(). Può essere utilizzato da qualsiasi altro thread CUDA anche da lanci successivi del kernel. Qualsiasi thread CUDA può liberare memoria allocata da un altro thread, ma è necessario prestare attenzione per garantire che lo stesso puntatore non venga liberato più di una volta.

Answer 1

Secondo http://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/CUDA_C_Programming_Guide.pdf dovresti poter usare malloc() e free() in una funzione del dispositivo.

Page 122

B.15 Dynamic Global allocazione memoria void * malloc (size_t size); libero da void (void * ptr); allocare e liberare memoria dinamicamente da un heap di dimensioni fisse nella memoria globale.

L'esempio riportato nel manuale.

__global__ void mallocTest() 
{ 
    char* ptr = (char*)malloc(123); 
    printf(“Thread %d got pointer: %p\n”, threadIdx.x, ptr); 
    free(ptr); 
} 

void main() 
{ 
    // Set a heap size of 128 megabytes. Note that this must 
    // be done before any kernel is launched. 
    cudaThreadSetLimit(cudaLimitMallocHeapSize, 128*1024*1024); 
    mallocTest<<<1, 5>>>(); 
    cudaThreadSynchronize(); 
} 

È necessario il parametro del compilatore -arch = sm_20 e una scheda che supporta> 2x architettura.