CUDA przydziela pamięć w funkcji __device__

Question

Czy istnieje sposób w CUDA do alokacji pamięci w funkcji __device__? Nie mogłem znaleźć żadnych przykładów takiego działania.

Od instrukcji: B.15 Dynamiczna globalna alokacja pamięci void * malloc (size_t size); void free (void * ptr); dynamicznie przydzielaj i zwalniaj pamięć ze sterty o stałym rozmiarze w pamięci globalnej. Funkcja malloc() jądra CUDA przydziela co najmniej bajty wielkości ze sterty urządzeń i zwraca wskaźnik do przydzielonej pamięci lub NULL, jeśli brakuje pamięci do spełnienia żądania. Zwrócony wskaźnik gwarantuje wyrównanie do 16-bajtowej granicy. Funkcja CUDA in-kernel free() zwolni pamięć wskazaną przez ptr, która musiała zostać zwrócona przez poprzednie wywołanie malloc(). Jeśli ptr ma wartość NULL, wywołanie funkcji free() jest ignorowane. Wielokrotne wywołania funkcji free() z tym samym ptr mają niezdefiniowane zachowanie. Pamięć przydzielona przez dany wątek CUDA za pomocą funkcji malloc() pozostaje przydzielona na cały okres istnienia kontekstu CUDA lub do momentu, gdy zostanie wyraźnie zwolniona przez wywołanie funkcji free(). Może być używany przez dowolne inne wątki CUDA nawet z kolejnych uruchomień jądra. Każdy wątek CUDA może zwolnić pamięć przydzieloną przez inny wątek, ale należy zadbać o to, aby ten sam wskaźnik nie był zwalniany więcej niż jeden raz.

Answer 1

Zgodnie z http://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/CUDA_C_Programming_Guide.pdf powinieneś móc używać funkcji malloc() i free() w funkcji urządzenia.

Page 122

B.15 Dynamiczny Global Allocation pamięci void * malloc (size_t size); void free (void * ptr); dynamicznie przydzielaj i zwalniaj pamięć ze sterty o stałym rozmiarze w pamięci globalnej.

Przykład podany w instrukcji.

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

Potrzebny jest parametr kompilatora -arch = sm_20 i karta obsługująca> 2x architektura.