CUDA: in quale spazio di memoria è memorizzata una matrice di dimensioni fisse?

Question

Quando si imposta una matrice di dimensione fissa in un kernel, ad esempio:

int my_array[100]; 

in cui lo spazio di memoria non dell'array finire?

In particolare, vorrei sapere se tale array può essere memorizzato nel file di registro o nella memoria condivisa su dispositivi> = 2.0 e, in tal caso, quali sono i requisiti.

Answer 1

Per Fermi (e architetture probabilmente precedenti), per un una matrice da memorizzare nel file di registro, il seguente condizioni devono essere soddisfatte:

1. la matrice viene indicizzata solo con costanti
2. ci sono registri disponibili
3. Eventualmente, il compilatore fa anche alcune analisi per determinare l'impatto sulle prestazioni complessive

Il motivo per (1) è che gli indici di registro sono codificati direttamente all'interno delle istruzioni SASS. Non c'è modo di indirizzare i registri indirettamente.

I principali fattori che limita il numero di registri per (2) sono: istruzioni

• Sass contengono solo 6 bit per il registro indicizzazione, che limita il numero di registri che può essere utilizzato in un kernel per 64. Il numero effettivo è 63 quindi uno è riservato a qualcosa.
• Un SM ha un blocco di registri che sono condivisi da tutti i thread che sono contemporaneamente in volo.
• I registri sono anche necessari per contenere le variabili, quindi il compilatore deve bilanciare l'utilizzo del registro per ottenere le migliori prestazioni generali.

Un potenziale rimedio per (1) è lo srotolamento del ciclo. Se un loop utilizza un contatore di loop come indice in un array, srotolare il loop (con #pragma unroll o manualmente) fa sì che gli indici di array diventino costanti in quanto vi è ora un'istruzione SASS separata per ciascun accesso di array.

Basato in parte su questa presentazione NVIDIA: Local Memory and Register Spilling. Il documento entra anche nei dettagli su come le posizioni di variabili e array influenzano le prestazioni.

Answer 2

Gli array locali all'interno di un kernel, come quelli definiti sono allocati nei registri e nella memoria locale quando il registro non è sufficiente.

Se si desidera allocare la matrice nella memoria condivisa è necessario specificare come segue:

__shared__ int my_array[100];