allocazione di memoria all'interno di un kernel CUDA

Question

Ho il seguente (snippet) di un kernel.

__global__ void plain(int* geneVec, float* probs, int* nComponents, float* randomNumbers,int *nGenes) 
{ 

    int xid = threadIdx.x + (blockDim.x * blockIdx.x); 

    float* currentProbs= (float*)malloc(sizeof(float)*tmp); 

     ..... 
     ..... 

    currentProbs[0] = probs[start]; 
    for (k=1;k<nComponents[0]; k++) 
    { 
     currentProbs[k] = currentProbs[k-1] + prob; 
    } 

     ... 
     ... 
     free(currentProbs); 

} 

Quando è statici (anche le stesse dimensioni) è molto veloce, ma quando CurrentProbs è allocato dinamicamente (come sopra) prestazioni è pessimo.

Questa domanda ha detto che potevo fare questo all'interno di un kernel: CUDA allocate memory in __device__ function

Ecco una questione connessa: Efficiency of Malloc function in CUDA

Mi chiedevo se altri metodi hanno risolto questo diverso da quello proposto nella carta? Sembra ridicolo che non si possa malloc/free all'interno di un kernel senza questo tipo di penalità.

Answer 1

Penso che la ragione per cui l'introduzione di malloc() rallenta il codice è che alloca la memoria nella memoria globale. Quando si utilizza una matrice di dimensioni fisse, è probabile che il compilatore la inserisca nel file di registro, che è molto più veloce.

Avere a che fare con un malloc all'interno del proprio kernel può significare che si sta provando a fare troppo lavoro con un singolo kernel. Se ogni thread assegna una quantità diversa di memoria, ogni thread esegue un numero diverso di volte nel ciclo for, e si ottiene molta divergenza di ordito.

Se ciascun thread in un warp esegue cicli ripetuti lo stesso numero di volte, è sufficiente allocarlo in primo piano. Anche se corrono un numero diverso di volte, puoi usare una dimensione costante. Invece, penso che dovresti considerare come puoi refactoring il tuo codice per rimuovere completamente quel loop dal tuo kernel.