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



Answers

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.

Question

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.




Related