PROBLEMA Tempistiche (ms) calcolo sommatoria di una riga di una matrice

[Tidusquall]

Buongiorno,
cerco di esporre la domanda nel modo più chiaro possibile.

La consegna dell'esercizio è la seguente:

Calcolare la sommatoria di una riga di una matrice con i seguenti parametri:
- N righe, con N = n°blocchi * n°thread
- X colonne
il risultato è un array di N elementi contenente le sommatorie di ogni riga

Il codice funziona (allego progetto), ma la mia perplessità è un altra: i tempi di esecuzione.
Per avere dei risultati più attendibili, eseguo 10000 cicli e calcolo la media dei tempi di esecuzione ma i risultati non mi convincono molto.

Come è possibile che tra i risultati mono e multi ci sia una differenza di quasi il 10% in più?

Le chiamate che faccio al kernel sono sia in mono che multi, ed i tempi risultano essere questi:

####################


MONO: 8 - 8 MULTI: 4x2 - 4x2

Time Mono: 0.117128 (ms) Time Multi: 0.094554 (ms)



MONO: 16 - 16 MULTI: 4x4 - 4x4

Time Mono: 0.206620 (ms) Time Multi: 0.193303 (ms)



MONO: 32 - 16 MULTI: 8x4 - 4x4

Time Mono: 0.220927 (ms) Time Multi: 0.204934 (ms)



MONO: 16 - 32 MULTI: 4x4 - 8x4

Time Mono: 0.300701 (ms) Time Multi: 0.274282 (ms)



MONO: 16 - 32 MULTI: 4x4 - 4x8

Time Mono: 0.300746 (ms) Time Multi: 0.274242 (ms)



MONO: 64 - 32 MULTI: 8x8 - 8x4

Time Mono: 0.562803 (ms) Time Multi: 0.453662 (ms)



MONO: 64 - 64 MULTI: 8x8 - 8x8

Time Mono: 1.346480 (ms) Time Multi: 0.874964 (ms)

Comprendo che ci debba essere una differenza tra mono e multi, ma addirittura nell'ultimo output siamo a 0.5 ms in più, non capisco.

 

[DispatchCode]

Ciao, ti ho editato il post, il codice pubblicalo pure sotto tag CODE la prossima volta.


In merito a CUDA, taggo @BrutPitt, che penso sappia risponderti.

 

[Tidusquall]

Ciao, ti ho editato il post, il codice pubblicalo pure sotto tag CODE la prossima volta.


In merito a CUDA, taggo @BrutPitt, che penso sappia risponderti.

Grazie mille per la dritta e per aver taggato un esperto!
Non so se il problema si limiti a CUDA o se generalmente il comportamento è uguale in qualsiasi ambito, ma sentire il parere di un esperto in materia è sicuramente utile.
Più che altro mi piacerebbe capire se è normale o meno e capire il motivo di tale differenza in ms (a patto che il codice sia corretto).

 

[BrutPitt]

Ringrazio @DispatchCode per la fiducia.
Intervengo per dire che purtroppo non saro' in sede (casa/studio) fino a Venerdi', quindi sono senza i miei "strumenti" (PC con cuda) per poter valutare il codice in esecuzione... e le mie competenze in CUDA non sono alla stregua di altri linguaggi di GPU Compute (OpenGL/Vulkan)

Ma ad una prima "rapida" analisi (ed esclusivamente visiva: sono in pausa pranzo presso un'azienda) mi verrebbe da dire che:

for(int j = 0; j<CICLI; ++j)
    {
        cudaEventCreate(&start);
        cudaEventCreate(&stop);   // ??? ... Deve essere prima dell'esecuzione
        cudaEventRecord(start, NULL);


        // Invoke kernel multi
        MatSumMultiKernel<<<dimGrid, dimBlock>>>(d_A, d_S);
    
        // stop and destroy timer
        cudaEventRecord(stop, NULL);
        cudaEventSynchronize(stop);
        cudaEventElapsedTime(&msecTotal, start, stop);

        //printf( "\n\nTime Multi: %f (ms)\n\n", msecTotal);
        mediat_multi += msecTotal;

        if(verbose)
        {
            cudaMemcpy(S_multi.elements, d_S.elements, size, cudaMemcpyDeviceToHost);
        }
    }

Ma avrei bisogno di un po' piu' di tempo per valutare il codice, e ripeto, attualmente non ho modo di veridficare cio' che dico.
Spero stasera di poter fare di meglio.

 

[BrutPitt]

Intanto ti chiedevo un paio di cose (che non ho modo di verificare), perche' ho la sensazione che vi siano delle inconguenze:

Stai parlando dei tempi di esecuzione, ma i dati sono congruenti (che e' la cosa di primaria importanza)?
Hai verificato che i due vettori siano identici, anche con l'utilizzo di un "terzo" vettore-somma di verifica, il cui calcolo venga effettuato via CPU, nel modo tradizionale (e senza parallelizzazione)?

Per cio' che riguarda cudaEventCreate(&stop) ... il risultato non cambiera' di molto ma almeno la chiamata e l'inizializzazione viene fatta "fuori" dal timer che e' gia' in esecuzione.

 

[Tidusquall]

Intanto ti chiedevo un paio di cose (che non ho modo di verificare), perche' ho la sensazione che vi siano delle inconguenze:

Stai parlando dei tempi di esecuzione, ma i dati sono congruenti (che e' la cosa di primaria importanza)?
Hai verificato che i due vettori siano identici, anche con l'utilizzo di un "terzo" vettore-somma di verifica, il cui calcolo venga effettuato via CPU, nel modo tradizionale (e senza parallelizzazione)?

Per cio' che riguarda cudaEventCreate(&stop) ... il risultato non cambiera' di molto ma almeno la chiamata e l'inizializzazione viene fatta "fuori" dal timer che e' gia' in esecuzione.

Intanto grazie mille per il supporto!
Comunque sì, pare che i vettori siano identici, ho provato con degli output e li ho paragonati.
Il cudaEventCreate(&stop) effettivamente va messo prima, erroneamente era all'interno del timer che era già in esecuzione, grazie per avermelo fatto notare. Da quel che ho appurato comunque i dati mi sembrano congruenti.

Grazie ancora per l'aiuto! Fammi sapere se hai bisogno di altre info, nel frattempo provo a ragionarci su anche io.

 

[BrutPitt]

Purtroppo sono in trasferta fino a venerdi', ma mi sono fatto prestare un portatile da un collega con una NVidia 960, su cui ho potuto installare CUDA (grazie al lento WiFi dell'albergo: quasi 3 ore, per 4GB)

Intanto grazie mille per il supporto!
Comunque sì, pare che i vettori siano identici, ho provato con degli output e li ho paragonati.

Il cudaEventCreate(&stop) effettivamente va messo prima, erroneamente era all'interno del timer che era già in esecuzione, grazie per avermelo fatto notare. Da quel che ho appurato comunque i dati mi sembrano congruenti.

Grazie ancora per l'aiuto! Fammi sapere se hai bisogno di altre info, nel frattempo provo a ragionarci su anche io.

Si' ho controllato i valori e sono corretti, anche utilizzando un vettore di controllo calcolato sulla CPU

Per cio' che riguarda le esigue differenze, soprattutto sulle matrici piu' "piccole", credo che possa dipendere da questo:

tant'e' che sulle ultime matrici il divario aumenta, in percentuale.


I miei dubbi, oggi, derivavano dall'utilizzo del ciclo for.
Ora tu sei vincolato ai parametri dell'esercizio, ma in una sommatoria di questo tipo si dovrebbe evitare (per quanto possibile) il ciclo for.
Lo scopo di Cuda sarebbe la massima parallelizzazione, invece finche' non termina l'elaborazione di quel kernel (con il for), quel thread/kernel tiene impegnato un cudaCore... ed ogni kernel viene eseguito parallelamente sul singolo cudaCore.

Ti faccio un esempio di cio che vorrei dire: con il tuo kernel ed una matrice 16 righe x 100 colonne, useresti "solo" 16 cores in parallelo (kernel eseguiti su cudaCores) per la somma delle 100 colonne (che tengono impegnato il singolo cudaCore, col ciclo for), invece dei 1600 cudaCores (o quanti te ne concede la GPU) teoricamente utilizzabili contemporaneamente.

Si otterrebbe invece completa parallelizzazione, facendo qualcosa del genere:

__global__ void rowSum(const Matrix A, Sums S, unsigned maxSize)
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if(index<maxSize) {
        int x = index/A.height;
        int y = index%A.height;
        S.elements[y] += A.elements[y*A.width+x];        // ATTENZIONE!!! Somma senza sincronizzazione
    }
}
// con l'invocazione del tipo ...
{
       int threadsPerBlock = 64;
       int blocksPerGrid = (A.height*A.width + threadsPerBlock - 1) / threadsPerBlock;
       rowSum<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_S, A.height*A.width);
}

Pero' bisognerebbe poi integrarlo con i vincoli imposti dalla traccia del tuo esercizio.

P.S.
C'e' anche il problema che 2 (o piu') thread possano accedere contemporaneamente allo stesso elemento della somma, alterando il risultato.
Per questo motivo ed in questo caso, bisogna usare atomicAdd, al posto della somma diretta

atomicAdd(&S.elements[y],A.elements[y*A.width+x]);

Mentre nel tuo codice non serve in quanto il valore viene sommato su una variabile "locale" al kernel e non su una locazione di memoria condivisa.