[CUDA C] Instruzioni atomiche

[HRBF]

Ciao a tutti, mi sono appena iscritto anche se è da un po che leggo su questo forum.
Sto scrivendo una piccola applicazione in CUDA per Windows 7, l'ambiente di sviluppo è Visual C++ 2008. Ho ancora le idee un po confuse su CUDA infatti sto incontrando alcune difficoltà nella stesura del codice del Kernel: devo fare in modo che ciascun thread, appartenente ad un blocco, aggiunga un valore ad un valore precedente memorizzato in una variabile allocata nella device memory. Ovviamente l'operazione deve essere fatta sequenzialmente per permettere a un thread di leggere l'incremento fatto precedemente da un altro thread. Per questo motivo ho la necessità di utilizzare un'istruzione atomica. E qui arriva il bello

In un normale programma il codice è cosi:

for(i = 0; i < N; i++)
{
        n = n + punti[i].z * peso;
	d = d + peso;
}

if (n != 0.0) {
        r = n / d;
        g->peso = r * pow(livelli[l].scala, 2);
}
else g->peso = 0.0;

Non riesco a capire come posso scrivere quelle due righe di codice nel ciclo for per fare in modo che ogni thread possa sommare un valore al valore esistente.
Il mio tentativo è stato quello di usare la funzione "atomicAdd()" in questo modo:

codice kernel:

int i = blockDim.x * blockIdx.x + threadIdx.x;

atomicAdd(&n, punti[i].z * peso);
atomicAdd(&d, peso);

__syncthreads();

if (n != 0.0) {
	r = n / d;
	g->peso = r * pow(livelli[l].scala, 2);
}
else g->peso = 0.0;

ma a questo punto il compilatore mi da errore: "identifier "atomicAdd" is undefined"
Devo aggiungere qualche altra libreria??
Avete qualche altro suggerimento su come potrei fare eseguire quelle istruzioni?

Poi ho un altro dubbio sempre sul codice del kernel che ho riportato: come posso fare in modo che il codice dopo alla funzione "__syncthreads();" venga eseguito da 1 solo thread al termine di tutta la computazione??

L'hardware che ho a disposizione è una NVIDIA GT230M compute capability 1.2.

Grazie in anticipo!

[Tommo]

devi includere nel sorgente CUDA l'header che definisce le atomics, non mi ricordo qual'è.

Per far eseguire il codice ad un solo thread è facile, basta fare

if( i == n )

E lo eseguirà solo il thread n.

Purtroppo per te il bello arriva ora, perchè ti accorgerai che la GPU andrà più lenta della CPU
E probabilmente anche di diverse volte.

I thread della GPU presi singolarmente sono decisamente "deboli" (nell'hardware non sono nemmeno veri threads) e la serializzazione porta al completo sconvolgimento della pipeline.
Serializzare OGNI thread con TUTTI quelli precedenti porta al blocco completo.

L'unico modo che hai di andare veramente veloce è dimenticarti delle atomics e trovare un modo di fare la stessa cosa leggendo meno dati possibili dalla memoria, meno che meno dagli altri threads.
Così com'è ora fai meglio a lasciar stare.

[HRBF]

Innanzitutto grazie per la risposta!
Ho capito è meglio evitare di usare le funzioni atomiche, in effetti è un po un controsenso usare un'architettura per lavorare in parallelo e poi fargli eseguire delle operazioni in serie.
Praticamente mi suggerisci di fare eseguire quella parte di codice sulla CPU? Posso fare cosi, mi semplificherebbe molto la vita, il fatto è che quelle due istruzioni, in particolare nell'ultimo passo dell'algoritmo, devono essere eseguite circa 260 milioni di volte e non era male se riuscivo a passare il carico di lavoro sulla GPU.

Ho pensato a una soluzione alternativa di questo tipo:

int i = blockDim.x * blockIdx.x + threadIdx.x;

peso = .................
n[i] = punti[i].z * peso;
d[i] = peso;

__syncthreads();

Ogni thread di un blocco inserisce un valore in un vettore ("n" e "d") allocati nella shared memory. Quando tutti i thread hanno finito, eseguo la somma di tutti i valori del vettore sfruttando uno o più thread. Però in questo caso il problema è che i vettori "n" e "d" saranno visibili solo ai thread che appartengono ad un solo blocco....se io necessito di più blocchi per portare a termine l'operazione come posso fare?
Che ne pensate? Vi sembra una soluzione fattibile?

Qualsiasi idea o suggerimento è bene accetto

[Tommo]

Beh io ho detto il caso estremo, però in realtà operazioni seriali le puoi anche fare

Di solito si parte risolvendo tutto nella shared memory; poi si accumula il valore in un registro che viene riversato nella memoria globale. A quel punto un blocco copia tutti questi valori dalla memoria globale nella sua shared memory e li accumula a loro volta, finchè non rimane uno solo.
Sembra facile (o anche no), ma ci sono tantissime ottimizzazioni possibili, tipo sul numero di registri, sui bank conflicts, etc...

Prova a vedere i vari algoritmi di "gather", sul forum di CUDA ce ne sono tantissimi decisamente ottimizzati, che fanno un sacco di roba tra cui la somma di array che serve a te.

[HRBF]

Ciao!

Ho cercato sul forum di CUDA come mi hai suggerito e ho trovato molto materiale davvero interessante, in particolare gli algoritmi di riduzione che effettuano appunto le somme che servono a me. Sono riuscito ad applicarne uno con diverse ottimizzazioni.

Grazie mille per i suggerimenti! Credo che in futuro ritornerò a chiedere consigli

[marco.r]

Procedendo nel modo che proponi non ottieni nulla dei vantaggi della computazione su GPU, fai prima ad eseguire du CPU che almeno ti risparmi il passaggio dei valor dall'altra.
La reduce o fold di un vettore, in una architettura parallela, va fatto in modo totalmente diverso (e piu'complicato), perche' prevede una sorta di approccio divide et impera sul vettore da sommare e che permette di eseguire l'operazione (nel caso naive) in tempo log(n). Se cerchi su internet ci sono un paio di paper di nvidia a riguardo.

[HRBF]

Come ho già scritto sopra ho applicato un algoritmo di reduce basato su balanced trees che ha complessità O(n).

[marco.r]

Originariamente inviato da HRBF

Come ho già scritto sopra ho applicato un algoritmo di reduce basato su balanced trees che ha complessità O(n).

uhm, ma se ha complessita' O(n) che vantaggio hai ?

[Tommo]

La GPU ha la capacità di eseguire algoritmi O(n) in circa O(1)
Lo scopo è creare un algoritmo O(n) dove ogni elemento ha il suo thread...

[HRBF]

Esattamente... la complessità è come quella di un algoritmo sequenziale che gira su CPU, ma c'è il vantaggio della computazione in parallelo eseguita da "n" thread.

[HRBF]

Ciao a tutti, rispolvero questo thread per chiedere un aiuto: ho la necessità di usare l'istruzione atomica "atomicInc()" in un kernel.
Ho letto che il compilatore nvcc compila di default con il comando "-arch sm_10" cioè utilizza la compute capability 1.0. Per utilizzare l'istruzione atomica bisogna avere minimo compute capability 1.1 (la mia è 1.2), pertanto bisogna aggiungere un parametro al compilatore: "-arch sm_12" che indica di compilare usando compute capability 1.2.
Ho provato in mille modi ad aggiungere quel parametro ma il compilatore mi spara sempre fuori l'errore: " error: identifier "atomicInc" is undefined"
Vi chiedo come si fa ad aggiungere quel parametro al compilatore in VisualStudio 2008? Uso il compilatore nvcc.

Mi sono scaricato l'SDK CUDA e c'è un esempio molto interessante che si chiama "Simple Atomic Intrinsics" il quale dimostra come usare le istruzioni atomiche. Mi sono guardato il codice, riesco a compilarlo e eseguirlo e funziona perfettamente...... ma il bello è che ho guardato di quali include fa uso, ho guardato le proprietà del progetto.......ho usato le stesse include nel mio progetto e ho usato le stesse impostazioni come nell'esempio del SDK CUDA, ma l'esempio funziona, invece a me continua a spararmi l'errore sull'istruzione atomica non capisco perchè

Qualsiasi suggerimento è ben accetto!!

[gugoXX]

Originariamente inviato da Tommo

La GPU ha la capacità di eseguire algoritmi O(n) in circa O(1)
Lo scopo è creare un algoritmo O(n) dove ogni elemento ha il suo thread...

Caricare (load) la GPU con tutti gli input ha complessita' O(n) (li carichi serialmente)
Se quindi i calcoli sono banali e sono bloccanti, penso che convenga farli fare direttamente alla CPU.

[HRBF]

Grazie per la risposta, ma quello che avevo scritto sopra non c'entra niente con la mia ultima richiesta. Visto che si tratta sempre di istruzioni atomiche non volevo aprire un altro thread.

Comunque devo usare la "atomicInc()" solo in un blocco al termine dell'esecuzione di tutti gli altri blocchi......so benissimo che le istruzioni atomiche rallentano ma devo usarla perforza.......tu sapresti come fare per usare quella dannata atomicInc?

[Torav]

L'opzione da passare è -arch=sm_xx, dove nel tuo caso xx è 12 (se non ho capito male!) nvcc --help dovrebbe darti tutte le opzioni possibili

[HRBF]

Originariamente inviato da Torav

L'opzione da passare è -arch=sm_xx, dove nel tuo caso xx è 12 (se non ho capito male!) nvcc --help dovrebbe darti tutte le opzioni possibili

Grazie per la risposta, ho provato ma non compila mi dice: " fatal error C1189: #error : -- incorrect inclusion of a cudart header file c:\cuda\include\crt\func_macro.h "

Ho provato a passare al compilatore "-arch sm_12" oppure "-arch=sm_12" se non mi da l'errore di inclusione dell'header file che ho riportato qui sopra, mi dice sempre "atomicInc() undefined"

La cosa assurda è che ho fatto una copia del progetto dall'SDK CUDA (se lo compilo funziona) quello dove mostra un esempio di utilizzo delle funzioni atomiche, ci ho messo i miei sorgenti, e compilando mi da sempre l'errore sull'atomicInc()

Boh non so più cosa pensare.

[Torav]

Ma per gli include usi path relativi o metti solo #include <cuda_runtime.h> o che so io? Perché in questo caso devi dire al compilatore dove cercare i file header (e io non so come questo si faccia con visual studio).

[HRBF]

Originariamente inviato da Torav

Ma per gli include usi path relativi o metti solo #include <cuda_runtime.h> o che so io? Perché in questo caso devi dire al compilatore dove cercare i file header (e io non so come questo si faccia con visual studio).

Ciao, per gli include uso la forma classica: #include <cuda_runtime.h> gli ho passato anche i percorsi degli header al compilatore ma niente, sempre errore sulla funzione atomica.
Ci ho rinunciato, ho trovato un altro algoritmo che non fa uso della funzione atomica