performance guida - Come faccio a scegliere la griglia e le dimensioni dei blocchi per i kernel CUDA?





programmazione programming (4)


Ci sono due parti per quella risposta (l'ho scritta). Una parte è facile da quantificare, l'altra è più empirica.

Vincoli hardware:

Questa è la parte facile da quantificare. L'Appendice F dell'attuale guida alla programmazione CUDA elenca una serie di limiti rigidi che limitano il numero di thread per blocco che può essere lanciato da un kernel. Se superi uno di questi, il tuo kernel non funzionerà mai. Possono essere sintetizzati sommariamente come:

  1. Ogni blocco non può avere più di 512/1024 thread in totale ( capacità di calcolo 1.x o 2.xe successive)
  2. Le dimensioni massime di ogni blocco sono limitate a [512,512,64] / [1024,1024,64] (Calcolo 1.x / 2.xo successivo)
  3. Ogni blocco non può consumare più di 8k / 16k / 32k / 64k / 32k / 64k / 32k / 64k / 32k / 64k registri totali (Compute 1.0,1.1 / 1.2,1.3 / 2.x- / 3.0 / 3.2 / 3.5-5.2 / 5.3 / 6-6,1 / 6,2 / 7,0)
  4. Ogni blocco non può consumare più di 16kb / 48kb / 96kb di memoria condivisa (Compute 1.x / 2.x-6.2 / 7.0)

Se rimani entro questi limiti, qualsiasi kernel che riesci a compilare correttamente si avvierà senza errori.

Ottimizzazione delle prestazioni:

Questa è la parte empirica. Il numero di thread per blocco scelto all'interno dei vincoli hardware descritti sopra può influire sulle prestazioni del codice in esecuzione sull'hardware. Il modo in cui ogni codice si comporta sarà diverso e l'unico vero modo per quantificarlo è attraverso un attento benchmarking e profilazione. Ma di nuovo, molto sintetizzato in modo approssimativo:

  1. Il numero di thread per blocco dovrebbe essere un multiplo rotondo della dimensione del warp, che è 32 su tutto l'hardware corrente.
  2. Ciascuna unità multiprocessore in streaming sulla GPU deve disporre di una quantità sufficiente di warp attivi per nascondere sufficientemente tutta la diversa latenza dell'architettura e della memoria dell'architettura e ottenere il massimo throughput. L'approccio ortodosso qui è quello di cercare di ottenere una occupazione ottimale dell'hardware (a cui si riferisce la risposta di Roger Dahl ).

Il secondo punto è un argomento enorme che dubito che qualcuno proverà a coprirlo in una singola risposta StackOverflow. Ci sono persone che scrivono tesi di dottorato sull'analisi quantitativa degli aspetti del problema (vedi questa presentazione di Vasily Volkov della UC Berkley e questo articolo di Henry Wong dell'Università di Toronto per esempi di quanto sia davvero complessa la domanda).

A livello di entrata, dovresti principalmente sapere che la dimensione del blocco che scegli (all'interno della gamma di dimensioni di blocco legali definite dai vincoli di cui sopra) può e ha un impatto sulla velocità con cui il tuo codice verrà eseguito, ma dipende dall'hardware hai e il codice che stai utilizzando. Con l'analisi comparativa, probabilmente troverai che la maggior parte del codice non banale ha un "punto debole" nella gamma di 128-512 thread per blocco, ma richiederà alcune analisi da parte tua per trovare dove si trova. La buona notizia è che, poiché si lavora in multipli della dimensione del warp, lo spazio di ricerca è molto limitato e la migliore configurazione per un dato pezzo di codice è relativamente facile da trovare.

Questa è una domanda su come determinare la griglia CUDA, il blocco e le dimensioni del filetto. Questa è una domanda aggiuntiva a quella pubblicata qui:

https://.com/a/5643838/1292251

Seguendo questo link, la risposta di talonmie contiene un frammento di codice (vedi sotto). Non capisco il commento "valore solitamente scelto da tuning e vincoli hardware".

Non ho trovato una buona spiegazione o chiarimento che lo spieghi nella documentazione CUDA. In sintesi, la mia domanda è come determinare il blocco ottimale (= numero di thread) dato il seguente codice:

const int n = 128 * 1024;
int blocksize = 512; // value usually chosen by tuning and hardware constraints
int nblocks = n / nthreads; // value determine by block size and total work
madd<<<nblocks,blocksize>>>mAdd(A,B,C,n);

A proposito, ho iniziato la mia domanda con il link sopra perché in parte risponde alla mia prima domanda. Se questo non è il modo corretto di porre domande su , si prega di scusarmi o di consigliarmi.




Le risposte sopra riportate indicano come la dimensione del blocco può influire sulle prestazioni e suggerisce un'euristica comune per la sua scelta basata sulla massimizzazione dell'occupazione. Senza voler fornire il criterio per scegliere la dimensione del blocco, vale la pena ricordare che CUDA 6.5 (ora nella versione Release Candidate) include diverse nuove funzioni di runtime per facilitare i calcoli di occupazione e avviare la configurazione, vedere

CUDA Pro Tip: Occupancy API semplifica la configurazione di avvio

Una delle funzioni utili è cudaOccupancyMaxPotentialBlockSize che calcola euristicamente una dimensione del blocco che raggiunge la massima occupazione. I valori forniti da tale funzione potrebbero quindi essere utilizzati come punto di partenza per l'ottimizzazione manuale dei parametri di avvio. Di seguito è un piccolo esempio.

#include <stdio.h>

/************************/
/* TEST KERNEL FUNCTION */
/************************/
__global__ void MyKernel(int *a, int *b, int *c, int N) 
{ 
    int idx = threadIdx.x + blockIdx.x * blockDim.x; 

    if (idx < N) { c[idx] = a[idx] + b[idx]; } 
} 

/********/
/* MAIN */
/********/
void main() 
{ 
    const int N = 1000000;

    int blockSize;      // The launch configurator returned block size 
    int minGridSize;    // The minimum grid size needed to achieve the maximum occupancy for a full device launch 
    int gridSize;       // The actual grid size needed, based on input size 

    int* h_vec1 = (int*) malloc(N*sizeof(int));
    int* h_vec2 = (int*) malloc(N*sizeof(int));
    int* h_vec3 = (int*) malloc(N*sizeof(int));
    int* h_vec4 = (int*) malloc(N*sizeof(int));

    int* d_vec1; cudaMalloc((void**)&d_vec1, N*sizeof(int));
    int* d_vec2; cudaMalloc((void**)&d_vec2, N*sizeof(int));
    int* d_vec3; cudaMalloc((void**)&d_vec3, N*sizeof(int));

    for (int i=0; i<N; i++) {
        h_vec1[i] = 10;
        h_vec2[i] = 20;
        h_vec4[i] = h_vec1[i] + h_vec2[i];
    }

    cudaMemcpy(d_vec1, h_vec1, N*sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_vec2, h_vec2, N*sizeof(int), cudaMemcpyHostToDevice);

    float time;
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start, 0);

    cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, MyKernel, 0, N); 

    // Round up according to array size 
    gridSize = (N + blockSize - 1) / blockSize; 

    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Occupancy calculator elapsed time:  %3.3f ms \n", time);

    cudaEventRecord(start, 0);

    MyKernel<<<gridSize, blockSize>>>(d_vec1, d_vec2, d_vec3, N); 

    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Kernel elapsed time:  %3.3f ms \n", time);

    printf("Blocksize %i\n", blockSize);

    cudaMemcpy(h_vec3, d_vec3, N*sizeof(int), cudaMemcpyDeviceToHost);

    for (int i=0; i<N; i++) {
        if (h_vec3[i] != h_vec4[i]) { printf("Error at i = %i! Host = %i; Device = %i\n", i, h_vec4[i], h_vec3[i]); return; };
    }

    printf("Test passed\n");

}

MODIFICARE

cudaOccupancyMaxPotentialBlockSize è definito nel file cuda_runtime.h ed è definito come segue:

template<class T>
__inline__ __host__ CUDART_DEVICE cudaError_t cudaOccupancyMaxPotentialBlockSize(
    int    *minGridSize,
    int    *blockSize,
    T       func,
    size_t  dynamicSMemSize = 0,
    int     blockSizeLimit = 0)
{
    return cudaOccupancyMaxPotentialBlockSizeVariableSMem(minGridSize, blockSize, func, __cudaOccupancyB2DHelper(dynamicSMemSize), blockSizeLimit);
}

Il significato per i parametri è il seguente

minGridSize     = Suggested min grid size to achieve a full machine launch.
blockSize       = Suggested block size to achieve maximum occupancy.
func            = Kernel function.
dynamicSMemSize = Size of dynamically allocated shared memory. Of course, it is known at runtime before any kernel launch. The size of the statically allocated shared memory is not needed as it is inferred by the properties of func.
blockSizeLimit  = Maximum size for each block. In the case of 1D kernels, it can coincide with the number of input elements.

Si noti che, a partire da CUDA 6.5, è necessario calcolare le proprie dimensioni di blocco 2D / 3D dalla dimensione del blocco 1D suggerita dall'API.

Si noti inoltre che l'API del driver CUDA contiene API funzionalmente equivalenti per il calcolo dell'occupazione, quindi è possibile utilizzare cuOccupancyMaxPotentialBlockSize nel codice dell'API del driver nello stesso modo mostrato per l'API di runtime nell'esempio sopra.




Il blocco è solitamente selezionato per massimizzare la "occupazione". Cerca su Occupazione CUDA per ulteriori informazioni. In particolare, consultare il foglio di calcolo CUDA Occupancy Calculator.








performance optimization cuda gpu nvidia