2012-11-09 3 views
40

Sto lavorando a un'applicazione statistica contenente circa 10-30 milioni di valori in virgola mobile in un array.Posso/devo eseguire questo codice su una GPU?

Diversi metodi di esecuzione di calcoli diversi, ma indipendenti, sulla matrice di cicli annidati, ad esempio:

Dictionary<float, int> noOfNumbers = new Dictionary<float, int>(); 

for (float x = 0f; x < 100f; x += 0.0001f) { 
    int noOfOccurrences = 0; 

    foreach (float y in largeFloatingPointArray) { 
     if (x == y) { 
      noOfOccurrences++; 
     } 
    } 

    noOfNumbers.Add(x, noOfOccurrences); 
} 

L'applicazione corrente è scritto in C#, gira su una CPU Intel e le esigenze diverse ore per completare. Non conosco i concetti e le API di programmazione GPU, quindi le mie domande sono:

  • È possibile (e ha senso) utilizzare una GPU per accelerare tali calcoli?
  • Se sì: qualcuno conosce un tutorial o ha qualche codice di esempio (il linguaggio di programmazione non è importante)?

Qualsiasi aiuto sarebbe molto apprezzato.

+2

Per caso, hai provato a convertire il codice in C/C++? In base allo snippet di codice riportato di seguito, stai utilizzando C#. Non sarei sorpreso se il tuo codice impiegasse molto tempo ad allocare memoria per il dizionario. – Martin

+3

No, ma l'allocazione della memoria per il dizionario richiede solo pochi ms o meno e l'utilizzo della CPU è sempre compreso tra 93% - 98%, quindi penso che la memoria non sia il problema di prestazioni (principale) in questo caso. – Mike

+4

Penso davvero che il tuo codice dovrebbe essere velocissimo senza usare una GPU. Hai provato ad allontanarti dall'uso di un dizionario (preallocato di tutto). Non usare foreach ma per. GPU è eccessivo. Riscrivi tutto in C, ti costringerà a pensare all'assegnazione della memoria. – Martin

risposta

76

UPDATE GPU Versione

__global__ void hash (float *largeFloatingPointArray,int largeFloatingPointArraySize, int *dictionary, int size, int num_blocks) 
{ 
    int x = (threadIdx.x + blockIdx.x * blockDim.x); // Each thread of each block will 
    float y;           // compute one (or more) floats 
    int noOfOccurrences = 0; 
    int a; 

    while(x < size)   // While there is work to do each thread will: 
    { 
     dictionary[x] = 0;  // Initialize the position in each it will work 
     noOfOccurrences = 0;  

     for(int j = 0 ;j < largeFloatingPointArraySize; j ++) // Search for floats 
     {              // that are equal 
                  // to it assign float 
      y = largeFloatingPointArray[j]; // Take a candidate from the floats array 
      y *= 10000;      // e.g if y = 0.0001f; 
      a = y + 0.5;      // a = 1 + 0.5 = 1; 
      if (a == x) noOfOccurrences++;  
     }          

     dictionary[x] += noOfOccurrences; // Update in the dictionary 
              // the number of times that the float appears 

    x += blockDim.x * gridDim.x; // Update the position here the thread will work 
    } 
} 

Questo quello che ho appena provato per gli ingressi più piccoli, perché sto testando ho il mio computer portatile. Tuttavia, ha funzionato. Tuttavia, è necessario fare ulteriori testicoli.

UPDATE sequenziale Versione

Ho appena fatto questa versione ingenua che esegue l'algoritmo per 30 milioni in meno di 20 secondi (funzione già conteggio per generare dati).

Fondamentalmente, ordina la tua gamma di galleggianti. Percorrerà la matrice ordinata, analizzando il numero di volte in cui un valore appare consecutivamente nell'array e quindi inserirà questo valore in un dizionario insieme al numero di volte in cui appare.

È possibile utilizzare la mappa ordinata, anziché la unordered_map che ho utilizzato.

Heres il codice:

#include <stdio.h> 
#include <stdlib.h> 
#include "cuda.h" 
#include <algorithm> 
#include <string> 
#include <iostream> 
#include <tr1/unordered_map> 


typedef std::tr1::unordered_map<float, int> Mymap; 


void generator(float *data, long int size) 
{ 
    float LO = 0.0; 
    float HI = 100.0; 

    for(long int i = 0; i < size; i++) 
     data[i] = LO + (float)rand()/((float)RAND_MAX/(HI-LO)); 
} 

void print_array(float *data, long int size) 
{ 

    for(long int i = 2; i < size; i++) 
     printf("%f\n",data[i]); 

} 

std::tr1::unordered_map<float, int> fill_dict(float *data, int size) 
{ 
    float previous = data[0]; 
    int count = 1; 
    std::tr1::unordered_map<float, int> dict; 

    for(long int i = 1; i < size; i++) 
    { 
     if(previous == data[i]) 
      count++; 
     else 
     { 
      dict.insert(Mymap::value_type(previous,count)); 
      previous = data[i]; 
      count = 1;   
     } 

    } 
    dict.insert(Mymap::value_type(previous,count)); // add the last member 
    return dict; 

} 

void printMAP(std::tr1::unordered_map<float, int> dict) 
{ 
    for(std::tr1::unordered_map<float, int>::iterator i = dict.begin(); i != dict.end(); i++) 
    { 
    std::cout << "key(string): " << i->first << ", value(int): " << i->second << std::endl; 
    } 
} 


int main(int argc, char** argv) 
{ 
    int size = 1000000; 
    if(argc > 1) size = atoi(argv[1]); 
    printf("Size = %d",size); 

    float data[size]; 
    using namespace __gnu_cxx; 

    std::tr1::unordered_map<float, int> dict; 

    generator(data,size); 

    sort(data, data + size); 
    dict = fill_dict(data,size); 

    return 0; 
} 

Se avete la spinta Libreria installata in te macchina che si dovrebbe usare questo:

#include <thrust/sort.h> 
thrust::sort(data, data + size); 

invece di questo

sort(data, data + size); 

Di sicuro sarà più veloce.

originale Messaggio

"Sto lavorando su un'applicazione statistica che ha una vasta gamma contenent 10 - 30 milioni di valori in virgola mobile".

"È possibile (e ha senso) utilizzare una GPU per accelerare tali calcoli?"

Sì, lo è. Un mese fa ho inserito una simulazione Molecular Dynamic interamente sulla GPU. Uno dei kernel, che calcola la forza tra coppie di particelle, riceve 6 array ciascuno con 500.000 doppi, per un totale di 3 milioni di doppi (22 MB).

Così si sta pianificando di inserire 30 milioni di punti float, si tratta di circa 114 MB di memoria globale, quindi questo non è un problema, anche il mio portatile ha 250 MB.

Il numero di calcoli può essere un problema nel tuo caso? Sulla base della mia esperienza con la Molecular Dynamic (MD), dico di no. La versione sequenziale di MD richiede circa 25 ore per essere completata mentre in GPU sono necessari 45 minuti. Hai detto che la tua applicazione impiega un paio d'ore, anche in base all'esempio di codice che sembra più morbido rispetto alla Dinamica molecolare.

Ecco l'esempio di calcolo forza:

__global__ void add(double *fx, double *fy, double *fz, 
        double *x, double *y, double *z,...){ 

    int pos = (threadIdx.x + blockIdx.x * blockDim.x); 

    ... 

    while(pos < particles) 
    { 

     for (i = 0; i < particles; i++) 
     { 
       if(//inside of the same radius) 
       { 
       // calculate force 
       } 
     } 
    pos += blockDim.x * gridDim.x; 
    }   
    } 

Un semplice esempio di codice nel Cuda potrebbe essere la somma di due matrici 2D:

In c:

for(int i = 0; i < N; i++) 
    c[i] = a[i] + b[i]; 

In Cuda :

__global__ add(int *c, int *a, int*b, int N) 
{ 
    int pos = (threadIdx.x + blockIdx.x) 
    for(; i < N; pos +=blockDim.x) 
     c[pos] = a[pos] + b[pos]; 
} 

In Cuda fondamentalmente preso ciascuno per iterazione e dividere per ciascun filo,

1) threadIdx.x + blockIdx.x*blockDim.x; 

Ogni blocco avere un ID da 0 a N-1 (N il numero massimo di blocchi), e ciascun blocco hanno un numero X di fili con un id da 0 a X-1.

1) Fornisce l'iterazione che ogni thread calcolerà in base al suo id e al blocco id in cui si trova il thread, il blockDim.x è il numero di thread di un blocco.

Quindi, se si dispone di 2 blocchi ciascuna con 10 thread e un N = 40, il:

Thread 0 Block 0 will execute pos 0 
Thread 1 Block 0 will execute pos 1 
... 
Thread 9 Block 0 will execute pos 9 
Thread 0 Block 1 will execute pos 10 
.... 
Thread 9 Block 1 will execute pos 19 
Thread 0 Block 0 will execute pos 20 
... 
Thread 0 Block 1 will execute pos 30 
Thread 9 Block 1 will execute pos 39 

Guardando al tuo codice che ho fatto questa bozza di quello che potrebbe essere in CUDA:

__global__ hash (float *largeFloatingPointArray, int *dictionary) 
    // You can turn the dictionary in one array of int 
    // here each position will represent the float 
    // Since x = 0f; x < 100f; x += 0.0001f 
    // you can associate each x to different position 
    // in the dictionary: 

    // pos 0 have the same meaning as 0f; 
    // pos 1 means float 0.0001f 
    // pos 2 means float 0.0002f ect. 
    // Then you use the int of each position 
    // to count how many times that "float" had appeared 


    int x = blockIdx.x; // Each block will take a different x to work 
    float y; 

while(x < 1000000) // x < 100f (for incremental step of 0.0001f) 
{ 
    int noOfOccurrences = 0; 
    float z = converting_int_to_float(x); // This function will convert the x to the 
              // float like you use (x/0.0001) 

    // each thread of each block 
    // will takes the y from the array of largeFloatingPointArray 

    for(j = threadIdx.x; j < largeFloatingPointArraySize; j += blockDim.x) 
    { 
     y = largeFloatingPointArray[j]; 
     if (z == y) 
     { 
      noOfOccurrences++; 
     } 
    } 
    if(threadIdx.x == 0) // Thread master will update the values 
     atomicAdd(&dictionary[x], noOfOccurrences); 
    __syncthreads(); 
} 

Devi usare atomicAdd perché thread diversi da blocchi diversi possono scrivere/leggere noOfOccurrences allo stesso tempo, quindi devi essere sicuro dell'esclusione reciproca.

Questo è solo un approccio che è possibile fornire le iterazioni del ciclo esterno ai thread anziché ai blocchi.

Tutorial

Il Dr Dobbs Journal serie CUDA: Supercomputing for the masses da Rob Farmer è eccellente e copre quasi tutto nelle sue quattordici rate. Inizia anche piuttosto delicatamente ed è quindi abbastanza amichevole per i principianti.

e anothers:

Date un'occhiata sul l'ultimo elemento, si trovano molti link per saperne di CUDA.

OpenCL: OpenCL Tutorials | MacResearch

+13

Questo è il tipo di risposta che mi sono unito a SO per ... Kudos! – DarkWanderer

+6

Beh, cosa posso dire, questa è la migliore risposta che abbia mai avuto su SO. Sei un genio, grazie e: vielen Dank! :-) E OpenCL e AMD ATI, hai esperienza con questa combinazione, qual è la tua opinione? – Mike

+6

Grazie mille, ecco qui sicuramente risposte di qualità più elevata. Non ho mai provato OpenCL ad essere sincero, ho appena lavorato con cuda e dispositivi NVIDIA (ad esempio, Tesla C2050), perché è disponibile nel cluster che ho ottenuto nel mio lavoro :). – dreamcrash

11

Non conosco molto dell'elaborazione parallela o GPGPU, ma per questo specifico esempio, è possibile risparmiare un sacco di tempo effettuando un singolo passaggio sull'array di input anziché eseguirne il ciclo un milione di volte. Con set di dati di grandi dimensioni, di solito, se possibile, si desidera eseguire le operazioni in un'unica passata. Anche se stai eseguendo più calcoli indipendenti, se è sullo stesso set di dati potresti ottenere una maggiore velocità eseguendoli tutti nello stesso passaggio, poiché otterrai una migliore localizzazione di riferimento in questo modo. Ma potrebbe non valerne la pena per l'aumento della complessità del tuo codice.

Inoltre, in realtà non si desidera aggiungere una piccola quantità a un numero in virgola mobile ripetutamente come quello, l'errore di arrotondamento si sommerà e non si otterrà ciò che si intendeva. Ho aggiunto un'istruzione if al mio esempio di sotto per verificare se gli input corrispondono al modello di iterazione, ma omettalo se non ne hai effettivamente bisogno.

non so alcun C#, ma una singola implementazione passaggio del campione sarebbe simile a questa:

Dictionary<float, int> noOfNumbers = new Dictionary<float, int>(); 

foreach (float x in largeFloatingPointArray) 
{ 
    if (math.Truncate(x/0.0001f)*0.0001f == x) 
    { 
     if (noOfNumbers.ContainsKey(x)) 
      noOfNumbers.Add(x, noOfNumbers[x]+1); 
     else 
      noOfNumbers.Add(x, 1); 
    } 
} 

Spero che questo aiuti.

+4

È possibile migliorare il codice utilizzando TryGet anziché ContainsKey e quindi noOfNumbers [x]. Usando TryGet si salva una ricerca nel dizionario, che è O (1) ammortizzata (cioè non sempre O (1)) ed è una O (1) costosa dal momento che un dizionario è un tipo di dati piuttosto complesso. In ogni caso +1 –

+3

Grazie ad entrambi per il vostro aiuto. È molto apprezzato e i tuoi suggerimenti verranno aggiunti presto alla mia applicazione. Purtroppo ho quasi 100 altri metodi che ritengo non possano essere ottimizzati molto di più. Anche se riesco a velocizzare tali calcoli del 90% utilizzando le ottimizzazioni del codice, potrebbero essere necessarie diverse ore per completare una CPU veloce. – Mike

+3

Inviaci il metodo completo con un set di dati limitato (e il tuo benchmark). Questo ci darà la possibilità di aiutarti molto di più. Sulla base di quello che sto vedendo nel tuo codice fino ad ora, sono abbastanza sicuro di poter raddoppiare la velocità del codice prima ancora di iniziare a utilizzare la GPU. – Martin

6

Oltre al suggerimento del poster di cui sopra, utilizzare il TPL (libreria parallela delle attività) quando appropriato per l'esecuzione in parallelo su più core.

L'esempio precedente potrebbe utilizzare Parallel.Foreach e ConcurrentDictionary, ma una configurazione di riduzione della mappa più complessa in cui l'array è suddiviso in blocchi, generando un dizionario che sarebbe quindi ridotto a un singolo dizionario, otterrebbe risultati migliori.

Non so se tutti i tuoi calcoli siano mappati correttamente alle funzionalità della GPU, ma dovrai comunque utilizzare un algoritmo di riduzione della mappa per mappare i calcoli ai core della GPU e quindi ridurre i risultati parziali in un singolo risultato, quindi potresti farlo sulla CPU prima di passare a una piattaforma meno familiare.

+3

Grazie per i vostri suggerimenti. Sto già usando il TPL ma a un livello più alto. Ciò significa che la mia app chiama diversi metodi paralleli che sembrano funzionare bene (utilizzo della CPU superiore al 90%). – Mike

6

Non sono sicuro che l'utilizzo di GPU sarebbe una buona corrispondenza dato che i valori di 'largerFloatingPointArray' devono essere recuperati dalla memoria. La mia comprensione è che le GPU sono più adatte per i calcoli autonomi.

Penso che trasformare questa singola applicazione di processo in un'applicazione distribuita in esecuzione su molti sistemi e modificare l'algoritmo dovrebbe accelerare notevolmente le cose, a seconda di quanti sistemi sono disponibili.

È possibile utilizzare il classico approccio "divide et impera". L'approccio generale che prenderei è il seguente.

Utilizzare un sistema per eseguire il preprocesso di 'largeFloatingPointArray' in una tabella hash o un database. Questo sarebbe fatto in un unico passaggio. Utilizzerebbe il valore in virgola mobile come chiave e il numero di occorrenze nell'array come valore. Lo scenario peggiore è che ogni valore si verifica solo una volta, ma è improbabile. Se largeFloatingPointArray continua a cambiare ogni volta che viene eseguita l'applicazione, la tabella hash in memoria ha senso. Se è statico, la tabella può essere salvata in un database di valori-chiave come Berkeley DB. Chiamiamo questo sistema 'lookup'.

Su un altro sistema, chiamiamolo 'principale', creare blocchi di lavoro e 'spargere' gli elementi di lavoro su N sistemi e 'raccogliere' i risultati non appena diventano disponibili. E. un oggetto di lavoro potrebbe essere semplice come due numeri che indicano l'intervallo su cui un sistema dovrebbe funzionare. Quando un sistema completa il lavoro, invia un array di occorrenze ed è pronto a lavorare su un'altra porzione di lavoro.

Le prestazioni sono migliorate perché non continuiamo a ripetere su LargeFloatingPointArray. Se il sistema di ricerca diventa un collo di bottiglia, può essere replicato su tutti i sistemi necessari.

Con un numero sufficiente di sistemi che funzionano in parallelo, dovrebbe essere possibile ridurre i tempi di elaborazione fino a minuti.

Sto lavorando a un compilatore per la programmazione parallela in C per i sistemi basati su molti core, spesso definiti microserver, che sono o saranno costruiti utilizzando più moduli "system-on-a-chip" all'interno di un sistema. I fornitori di moduli ARM includono Calxeda, AMD, AMCC, ecc. Intel probabilmente avrà anche un'offerta simile.

Ho una versione del compilatore funzionante, che potrebbe essere utilizzata per tale applicazione. Il compilatore, basato su prototipi di funzione C, genera un codice di rete C che implementa il codice di comunicazione tra processi (IPC) tra i sistemi. Uno dei meccanismi IPC disponibili è socket/tcp/ip.

Se hai bisogno di aiuto nell'implementazione di una soluzione distribuita, sarei felice di discuterne con te.

Aggiunto Nov 16 2012.

ho pensato un po 'di più l'algoritmo e penso che questo dovrebbe farlo in un unico passaggio. È scritto in C e dovrebbe essere molto veloce rispetto a quello che hai.

/* 
* Convert the X range from 0f to 100f in steps of 0.0001f 
* into a range of integers 0 to 1 + (100 * 10000) to use as an 
* index into an array. 
*/ 

#define X_MAX   (1 + (100 * 10000)) 

/* 
* Number of floats in largeFloatingPointArray needs to be defined 
* below to be whatever your value is. 
*/ 

#define LARGE_ARRAY_MAX (1000) 

main() 
{ 
    int j, y, *noOfOccurances; 
    float *largeFloatingPointArray; 

    /* 
    * Allocate memory for largeFloatingPointArray and populate it. 
    */ 

    largeFloatingPointArray = (float *)malloc(LARGE_ARRAY_MAX * sizeof(float));  
    if (largeFloatingPointArray == 0) { 
     printf("out of memory\n"); 
     exit(1); 
    } 

    /* 
    * Allocate memory to hold noOfOccurances. The index/10000 is the 
    * the floating point number. The contents is the count. 
    * 
    * E.g. noOfOccurances[12345] = 20, means 1.2345f occurs 20 times 
    * in largeFloatingPointArray. 
    */ 

    noOfOccurances = (int *)calloc(X_MAX, sizeof(int)); 
    if (noOfOccurances == 0) { 
     printf("out of memory\n"); 
     exit(1); 
    } 

    for (j = 0; j < LARGE_ARRAY_MAX; j++) { 
     y = (int)(largeFloatingPointArray[j] * 10000); 
     if (y >= 0 && y <= X_MAX) { 
      noOfOccurances[y]++; 
     } 
    } 
} 
+3

il lavoro può essere suddiviso tra una rete di macchine in una seconda volta; ma IMHO per miglioramenti economici (e spesso enormi) che utilizzano la potenza della GPU è di gran lunga migliore. Per quanto riguarda il tuo framework, come si confronta con MPI? :) – Pragmateek

+0

Grazie per tutte le informazioni e il codice c. Forse ho trovato una buona soluzione per il mio problema: http://bit.ly/Ta4aSL [PDF] Sembra molto promettente ... cosa ne pensi? – Mike

+0

Mike, È un modo interessante per sfruttare DirectX senza essere legato a una particolare GPU. Stavo pensando agli effetti collaterali, se ce ne sono. Mentre DirectX viene utilizzato attivamente, c'è qualche impatto sulle altre applicazioni che mostrano la grafica sul display? Provare a riprodurre un video di YouTube Media Player o Windows con e senza la tua app in esecuzione e vedere se si nota un peggioramento della qualità del video in riproduzione. Inoltre, sai se in futuro potresti dover scalare le funzionalità della workstation? Dato che fa tutto parte dell'ambiente Windows, penso che valga la pena provare. –

8

E 'possibile (e ha senso) di utilizzare una GPU per accelerare tali calcoli?

  • Sicuramente SI, questo tipo di algoritmo è in genere il candidato ideale per massiccia dati parallelismo elaborazione, le GPU cosa sono così bravi a.

Se sì: Qualcuno sa qualsiasi tutorial o ha qualche esempio di codice (linguaggio di programmazione non importa)?

  • Quando si vuole andare nella direzione GPGPU si hanno due alternative: CUDA o OpenCL.

    CUDA è maturo con molti strumenti ma le GPU NVidia sono centrali.

    OpenCL è uno standard in esecuzione su GPU NVidia e AMD e CPU. Quindi dovresti davvero favorirlo.

  • Per esercitazione ha una serie eccellente su CodeProject da Rob Farber: http://www.codeproject.com/Articles/Rob-Farber#Articles

  • Per il vostro specifico caso d'uso c'è un sacco di campioni per istogrammi buiding con OpenCL (da notare che molti sono gli istogrammi di immagini ma i principi sono gli stessi).

  • Come si usa C# è possibile utilizzare attacchi come OpenCL.Net o Cloo.

  • Se la matrice è troppo grande per essere archiviata nella memoria della GPU, è possibile suddividerla partizionalmente e rieseguire il kernel OpenCL per ciascuna parte facilmente.

+2

un'ulteriore risorsa su algos istogramma efficiente ... http://users.cecs.anu.edu.au/~ramtin/cuda.htm – kineticfocus

+2

Grazie per il vostro aiuto! Molto apprezzato. Qual è la tua opinione su DirectX? Sembra esserci un buon SDK per C# www.sharpdx.org – Mike

+2

Realizzato qualche ricerca aggiuntiva. OpenCL è molto interessante perché supporta anche Xeon Phi e la GPU integrata delle moderne CPU Intel, vedi qui http://bit.ly/Ta29ab – Mike