optimization reduction - Efficienza delle dimensioni di blocchi e griglie CUDA





kernel (5)


Se si dispone di set di dati dimensionati dinamicamente, è probabile che si verifichino alcuni problemi con la latenza mentre alcuni thread e blocchi attendono il completamento degli altri.

Questo sito ha delle ottime euristiche. Alcuni punti salienti generali:

Scelta dei blocchi per griglia

  • I blocchi per griglia dovrebbero essere> = numero di multiprocessori.
  • __syncthreads() uso di __syncthreads() nei kernel, più blocchi (in modo che un blocco possa essere eseguito mentre un altro attende la sincronizzazione)

Scelta di thread per blocco

  • Discussioni in multipli di dimensione di curvatura (cioè in genere 32)

  • Generalmente utile scegliere il numero di thread in modo che il numero massimo di thread per blocco (basato su hardware) sia un multiplo del numero di thread. Ad esempio con thread max di 768, l'utilizzo di 256 thread per blocco tenderà a essere migliore di 512 perché più thread possono essere eseguiti contemporaneamente su un blocco.

Qual è il modo consigliato di trattare i dataset di dimensioni dinamiche in cuda?

Si tratta di "impostare il blocco e le dimensioni della griglia in base al set di problemi" o vale la pena assegnare le dimensioni dei blocchi come fattori di 2 e disporre di una logica in-kernel per gestire l'over-spill?

Posso vedere come questo probabilmente importi molto per le dimensioni del blocco, ma quanto conta questo per le dimensioni della griglia? A mio avviso, i vincoli hardware effettivi si fermano a livello di blocco (cioè i blocchi assegnati a SM che hanno un determinato numero di SP, e quindi in grado di gestire una particolare dimensione di curvatura).

Ho esaminato la "Programmazione di processori paralleli in parallelo" di Kirk, ma in realtà non tocca proprio quest'area.




Ok suppongo che abbiamo a che fare con due domande qui.

1) Un buon modo per assegnare dimensioni del blocco (cioè il numero di thread) Questo di solito dipende dal tipo di dati che stai trattando. Hai a che fare con i vettori? Hai a che fare con le matrici? Il modo suggerito è di mantenere il numero di thread in multipli di 32. Quindi, quando si tratta di vettori, l'avvio di 256 x 1, 512 x 1 blocchi può andare bene. E similariy quando si tratta di matrici, 32 x 8, 32 x 16.

2) Un buon modo per assegnare le dimensioni della griglia (ovvero il numero di blocchi) Qui diventa un po 'complicato. È sufficiente lanciare 10.000 blocchi perché non è normalmente il modo migliore di fare le cose. Cambiare i blocchi dentro e fuori l'hardware è costoso. Due cose da considerare sono la memoria condivisa utilizzata per blocco e il numero totale di SP disponibili e risolvere il numero ottimale.

Puoi trovare una buona implementazione di come farlo dalla spinta . Tuttavia, potrebbe volerci un po 'per capire cosa sta succedendo all'interno del codice.




Solitamente si tratta di impostare la dimensione del blocco per prestazioni ottimali e la dimensione della griglia in base alla quantità totale di lavoro. La maggior parte dei kernel ha un numero "dolce" di warp per MP dove funzionano meglio, e dovresti fare un po 'di benchmarking / profiling per vedere dove sia. Probabilmente hai ancora bisogno di una logica di overflow nel kernel perché le dimensioni dei problemi sono raramente multipli rotondi di dimensioni dei blocchi.

EDIT: per dare un esempio concreto di come ciò potrebbe essere fatto per un kernel semplice (in questo caso un'operazione di tipo dscal BLAS di livello 1 eseguita come parte di una fattorizzazione di Cholesky di matrici di bande simmetriche compresse):

// Fused square root and dscal operation
__global__ 
void cdivkernel(const int n, double *a)
{
    __shared__ double oneondiagv;

    int imin = threadIdx.x + blockDim.x * blockIdx.x;
    int istride = blockDim.x * gridDim.x;

    if (threadIdx.x == 0) {
        oneondiagv = rsqrt( a[0] );
    }
    __syncthreads();

    for(int i=imin; i<n; i+=istride) {
        a[i] *= oneondiagv;
    }
}

Per avviare questo kernel, i parametri di esecuzione sono calcolati come segue:

  1. Consentiamo fino a 4 warps per blocco (quindi 128 thread). Normalmente si dovrebbe risolvere questo problema con un numero ottimale, ma in questo caso il kernel viene spesso chiamato su vettori molto piccoli, quindi avere una dimensione di blocco variabile ha un senso.
  2. Quindi calcoliamo il numero di blocchi in base alla quantità totale di lavoro, fino a 112 blocchi totali, che equivale a 8 blocchi per MP su un Fermi Telsa da 14 MP. Il kernel verrà iterato se la quantità di lavoro supera la dimensione della griglia.

La funzione wrapper risultante contenente i calcoli dei parametri di esecuzione e l'avvio del kernel hanno questo aspetto:

// Fused the diagonal element root and dscal operation into
// a single "cdiv" operation
void fusedDscal(const int n, double *a)
{
    // The semibandwidth (column length) determines
    // how many warps are required per column of the 
    // matrix.
    const int warpSize = 32;
    const int maxGridSize = 112; // this is 8 blocks per MP for a Telsa C2050

    int warpCount = (n / warpSize) + (((n % warpSize) == 0) ? 0 : 1);
    int warpPerBlock = max(1, min(4, warpCount));

    // For the cdiv kernel, the block size is allowed to grow to
    // four warps per block, and the block count becomes the warp count over four
    // or the GPU "fill" whichever is smaller
    int threadCount = warpSize * warpPerBlock;
    int blockCount = min( maxGridSize, max(1, warpCount/warpPerBlock) );
    dim3 BlockDim = dim3(threadCount, 1, 1);
    dim3 GridDim  = dim3(blockCount, 1, 1);

    cdivkernel<<< GridDim,BlockDim >>>(n,a);
    errchk( cudaPeekAtLastError() );
}

Forse questo fornisce alcuni suggerimenti su come progettare uno schema "universale" per l'impostazione dei parametri di esecuzione rispetto alle dimensioni dei dati di input.




Penso che di solito sia meglio impostare il blocco e le dimensioni della griglia in base al problema impostato, soprattutto per scopi di ottimizzazione. Avere thread aggiuntivi che non fanno nulla non ha senso e può peggiorare le prestazioni dei tuoi programmi.




"E. Se un warp contiene 20 thread, ma al momento sono disponibili solo 16 core, il warp non verrà eseguito."

non è corretto. Si stanno confondendo nuclei nel loro senso comune (usato anche nelle CPU) - il numero di "multiprocessori" in una GPU, con i nuclei nel marketing nVIDIA ("la nostra carta ha migliaia di core CUDA").

Un warp può essere programmato solo su un singolo core (= multiprocessore) e può eseguire fino a 32 thread contemporaneamente; non può usare più di un singolo core.

Il numero "48 warps" è il numero massimo di warps attivi (orditi che possono essere scelti per essere programmati per il lavoro nel ciclo successivo, in qualsiasi ciclo) per multiprocessore, su GPU nVIDIA con capacità di calcolo 2.x; e questo numero corrisponde a 1536 = 48 x 32 thread.

Risposta basata su questo webinar





optimization cuda gpgpu