CUDA asigna memoria en la función __device__

Question

¿Hay alguna manera en CUDA de asignar memoria en la función __device__? No pude encontrar ningún ejemplo de hacer esto.

De manual: B.15 Dynamic Global Memory Asignación void * malloc (size_t size); void free (void * ptr); asigna y libera memoria dinámicamente desde un montón de tamaño fijo en la memoria global. La función CUDA in-kernel malloc() asigna al menos un tamaño de bytes desde el montón del dispositivo y devuelve un puntero a la memoria asignada o NULL si no hay memoria suficiente para cumplir con la solicitud. Se garantiza que el puntero devuelto esté alineado con un límite de 16 bytes. La función CUDA in-kernel free() desasigna la memoria apuntada por ptr, que debe haber sido devuelta por una llamada previa a malloc(). Si ptr es NULL, la llamada a free() se ignora. Las llamadas repetidas a free() con el mismo ptr tienen un comportamiento indefinido. La memoria asignada por un hilo CUDA dado a través de malloc() permanece asignada durante el tiempo de vida del contexto CUDA, o hasta que se libera explícitamente mediante una llamada a free(). Puede ser utilizado por cualquier otro hilo CUDA, incluso desde el lanzamiento del núcleo posterior. Cualquier hilo CUDA puede liberar memoria asignada por otro hilo, pero se debe tener cuidado para asegurar que el mismo puntero no se libere más de una vez.

Answer 1

De acuerdo con http://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/CUDA_C_Programming_Guide.pdf, debería poder usar malloc() y free() en una función del dispositivo.

Página 122

B.15 dinámica global de asignación de memoria void * malloc (size_t tamaño); void free (void * ptr); asigna y libera memoria dinámicamente desde un montón de tamaño fijo en la memoria global.

El ejemplo dado en el manual.

__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(); 
} 

Se necesita el compilador parametro -arch = sm_20 y una tarjeta que soporta> 2x arquitectura.