cuda download - Memoria appuntata predefinita contro memoria a copia zero




toolkit nvidia (3)

In CUDA possiamo usare la memoria bloccata per copiare in modo più efficiente i dati da Host a GPU rispetto alla memoria predefinita allocata tramite malloc all'host. Tuttavia, esistono due tipi di memorie bloccate, la memoria bloccata predefinita e la memoria bloccata a zero copie .

La memoria bloccata predefinita copia i dati da Host a GPU due volte più veloci dei normali trasferimenti, quindi c'è sicuramente un vantaggio (a condizione che abbiamo abbastanza memoria host per bloccare la pagina)

Nella diversa versione della memoria bloccata, ovvero la memoria a zero copia , non è necessario copiare completamente i dati dall'host alla DRAM della GPU. I kernel leggono i dati direttamente dalla memoria dell'Host.

La mia domanda è: quale di questi tipi di memoria bloccata è una pratica di programmazione migliore.


Answers

Penso che dipenda dalla tua applicazione (altrimenti, perché dovrebbero fornire entrambi i modi?)

La memoria mappata e bloccata (copia zero) è utile quando:

  • La GPU non ha memoria da sola e utilizza comunque la RAM

  • Carichi i dati esattamente una volta, ma hai un sacco di calcoli da eseguire su di esso e vuoi nascondere le latenze di trasferimento della memoria attraverso di essa.

  • Il lato host vuole cambiare / aggiungere più dati, o leggere i risultati, mentre il kernel è ancora in esecuzione (es. Comunicazione)

  • I dati non si adattano alla memoria della GPU

Si noti che, è anche possibile utilizzare più flussi per copiare dati ed eseguire kernel in parallelo.

La memoria fissata, ma non mappata è migliore:

  • Quando carichi o memorizzi i dati più volte. Ad esempio: hai più kernel successivi, eseguendo il lavoro in fasi: non è necessario caricare i dati dall'host ogni volta.

  • Non c'è molto calcolo da eseguire e le latenze di caricamento non verranno nascoste bene


La memoria pinned mappata è identica ad altri tipi di memoria bloccata sotto tutti gli aspetti, tranne per il fatto che è mappata nello spazio degli indirizzi CUDA, quindi può essere letta e scritta dai kernel CUDA e utilizzata per i trasferimenti DMA dai Copy Engine.

Il vantaggio di non mappare la memoria bloccata era duplice: ti ha permesso di risparmiare spazio di indirizzamento, che può essere un bene prezioso in un mondo di piattaforme a 32 bit con GPU che possono contenere 3-4G di RAM. Inoltre, la memoria che non è mappata non può essere accidentalmente corrotta dai kernel rogue. Ma questa preoccupazione è abbastanza esoterica che la funzionalità di spazio di indirizzo unificata in CUDA 4.0 farà sì che tutte le assegnazioni a pin siano mappate per impostazione predefinita.

Oltre ai punti sollevati dal libro Sanders / Kandrot, altre cose da tenere a mente:

  • scrivere su host di memoria da un kernel (ad esempio per pubblicare i risultati sulla CPU) è bello in quanto la GPU non ha alcuna latenza da coprire in quel caso, e

  • È MOLTO IMPORTANTE che le operazioni di memoria siano coalizzate, altrimenti anche le GP 2.x e successive generano un grande successo di larghezza di banda.


Innanzitutto, ripeterò il mio commento: le GPU hanno una larghezza di banda elevata, una latenza elevata. Cercando di convincere la GPU a battere una CPU per un lavoro in nanosecondi (o anche un millisecondo o un secondo lavoro) manca completamente il senso di fare cose GPU. Di seguito è riportato un codice semplice, ma per apprezzare davvero i vantaggi prestazionali della GPU, sarà necessario un grande problema per ammortizzare i costi di avvio rispetto a ... altrimenti, è privo di significato. Posso battere una Ferrari in una corsa di due piedi, semplicemente perché ci vuole un po 'di tempo per girare la chiave, avviare il motore e premere il pedale. Ciò non significa che io sia più veloce della Ferrari in alcun modo significativo.

Usa qualcosa come questo in C ++:

  #define N (1024*1024)
  #define M (1000000)
  int main()
  {
     float data[N]; int count = 0;
     for(int i = 0; i < N; i++)
     {
        data[i] = 1.0f * i / N;
        for(int j = 0; j < M; j++)
        {
           data[i] = data[i] * data[i] - 0.25f;
        }
     }
     int sel;
     printf("Enter an index: ");
     scanf("%d", &sel);
     printf("data[%d] = %f\n", sel, data[sel]);
  }

Usa qualcosa di simile in CUDA / C:

  #define N (1024*1024)
  #define M (1000000)

  __global__ void cudakernel(float *buf)
  {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     buf[i] = 1.0f * i / N;
     for(int j = 0; j < M; j++)
        buf[i] = buf[i] * buf[i] - 0.25f;
  }

  int main()
  {
     float data[N]; int count = 0;
     float *d_data;
     cudaMalloc(&d_data, N * sizeof(float));
     cudakernel<<<N/256, 256>>>(d_data);
     cudaMemcpy(data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);
     cudaFree(d_data); 

     int sel;
     printf("Enter an index: ");
     scanf("%d", &sel);
     printf("data[%d] = %f\n", sel, data[sel]);
  }

Se ciò non funziona, prova a fare N e M più grande, o cambiare 256 a 128 o 512.