Sommatoria dell'array cummulativo usando OpenCL

Question

Sto calcolando la distanza euclidea tra i punti n-dimensionali usando OpenCL. Ottengo due liste di punti n-dimensionali e dovrei restituire un array che contenga solo le distanze da ogni punto della prima tabella a ogni punto della seconda tabella.

Il mio approccio è quello di fare il loop doble regolare (per ogni punto in Table1 {per ogni punto in Table2 {...}} e poi fare il calcolo per ogni coppia di punti in paralell.

euclidea la distanza viene quindi divisa in 3 parti: 1. prendere la differenza tra ogni dimensione nei punti 2. quadrare tale differenza (sempre per ogni dimensione) 3. sommare tutti i valori ottenuti in 2. 4. Prendere la radice quadrata del valore ottenuto in 3. (questo passaggio è stato omesso in questo esempio.)

Tutto funziona come un fascino finché io prova ad accumulare la somma di tutte le differenze (ovvero, eseguendo il passaggio 3. della procedura descritta sopra, riga 49 del codice seguente).

Come dati di test sto utilizzando le DescriptorList con 2 punti ciascuna: DescriptorList1: 001,002,003, ..., 127,128; (p1) 129,130,131, ..., 255,256; (P2)

DescriptorList2: 000.001.002, ..., 126.127; (p1) 128,129,130, ..., 254,255; (p2)

Quindi il vettore risultante dovrebbe avere i valori: 128, 2064512, 2130048, 128 Al momento sto ottenendo numeri casuali che variano ad ogni corsa.

Apprezzo qualsiasi aiuto o indizio su ciò che sto facendo male. Speriamo che tutto è chiaro sullo scenario sto lavorando in

#define BLOCK_SIZE 128 

typedef struct 
{ 
    //How large each point is 
    int length; 
    //How many points in every list 
    int num_elements; 
    //Pointer to the elements of the descriptor (stored as a raw array) 
    __global float *elements; 
} DescriptorList; 

__kernel void CompareDescriptors_deb(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float As[BLOCK_SIZE]) 
{ 

    int gpidA = get_global_id(0); 

    int featA = get_local_id(0); 

    //temporary array to store the difference between each dimension of 2 points 
    float dif_acum[BLOCK_SIZE]; 

    //counter to track the iterations of the inner loop 
    int loop = 0; 

    //loop over all descriptors in A 
    for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){ 

     //take the i-th descriptor. Returns a DescriptorList with just the i-th 
     //descriptor in DescriptorList A 
     DescriptorList tmpA = GetDescriptor(A, i); 

     //copy the current descriptor to local memory. 
     //returns one element of the only descriptor in DescriptorList tmpA 
     //and index featA 
     As[featA] = GetElement(tmpA, 0, featA); 
     //wait for all the threads to finish copying before continuing 
     barrier(CLK_LOCAL_MEM_FENCE); 

     //loop over all the descriptors in B 
     for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){ 
      //take the difference of both current points 
      dif_acum[featA] = As[featA]-B.elements[k*BLOCK_SIZE + featA]; 
      //wait again 
      barrier(CLK_LOCAL_MEM_FENCE); 
      //square value of the difference in dif_acum and store in C 
      //which is where the results should be stored at the end. 
      C[loop] = 0; 
      C[loop] += dif_acum[featA]*dif_acum[featA]; 
      loop += 1; 
      barrier(CLK_LOCAL_MEM_FENCE); 
     } 
    } 
} 

Answer 1

tuo problema risiede in queste righe di codice:.

C[loop] = 0; 
C[loop] += dif_acum[featA]*dif_acum[featA]; 

Tutte le discussioni del gruppo di lavoro (beh, in realtà tutte le discussioni, ma veniamo a quello più tardi) stiamo provando a modificare questa posizione dell'array contemporaneamente senza alcuna sincronizzazione. Diversi fattori rendono molto problematico:

1. Il gruppo di lavoro non è garantito il funzionamento completamente in parallelo, il che significa che per alcuni filati C [loop] = 0 può essere chiamato dopo altri thread hanno già eseguito la riga successiva
2. Quelli che eseguono in parallelo leggono tutti lo stesso valore da C [loop], lo modificano con il loro incremento e cercano di riscrivere lo stesso indirizzo. Non sono completamente sicuro di quale sia il risultato di questo writeback (penso che uno dei thread riesca a scrivere di nuovo, mentre gli altri falliscono, ma non ne sono completamente sicuro), ma è comunque sbagliato.

Ora lascia risolvere questo problema: Mentre potremmo essere in grado di ottenere questo al lavoro sulla memoria globale utilizzando Atomics, non sarà veloce, così lascia accumulare nella memoria locale:

local float* accum; 
... 
accum[featA] = dif_acum[featA]*dif_acum[featA]; 
barrier(CLK_LOCAL_MEM_FENCE); 
for(unsigned int i = 1; i < BLOCKSIZE; i *= 2) 
{ 
    if ((featA % (2*i)) == 0) 
     accum[featA] += accum[featA + i]; 
    barrier(CLK_LOCAL_MEM_FENCE); 
} 
if(featA == 0) 
    C[loop] = accum[0]; 

Di Certo, puoi riutilizzare altri buffer locali per questo, ma penso che il punto sia chiaro (btw: sei sicuro che dif_acum verrà creato nella memoria locale, perché penso di aver letto da qualche parte che questo non sarebbe stato messo nella memoria locale, che renderebbe preloading A nella memoria locale un po 'inutile).

Alcuni altri punti su questo codice:

1. Il codice è sembra essere orientato ad utilizzare solo il gruppo di lavoro (non si utilizza né groupid né ID Globale per vedere quali elementi su cui lavorare), per prestazioni ottimali che potresti voler usare di più.
2. potrebbe essere preferance personale, ma a me sembra meglio usare get_local_size(0) per la workgroupsize di utilizzare un define (poiché potrebbe cambiare nel codice ospitante senza rendersene conto avrebbe dovuto cambiare il codice OpenCL a)
3. Le barriere nel codice non sono necessarie, poiché nessun thread accede a un elemento nella memoria locale che è scritto da un altro thread. Pertanto non è necessario utilizzare la memoria locale per questo.

Considerando l'ultimo punto si potrebbe semplicemente fare:

float As = GetElement(tmpA, 0, featA); 
... 
float dif_acum = As-B.elements[k*BLOCK_SIZE + featA]; 

Ciò renderebbe il codice (non considerando i primi due proiettili):

__kernel void CompareDescriptors_deb(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float accum[BLOCK_SIZE]) 
{ 
    int gpidA = get_global_id(0); 
    int featA = get_local_id(0); 
    int loop = 0; 
    for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){ 
     DescriptorList tmpA = GetDescriptor(A, i); 
     float As = GetElement(tmpA, 0, featA); 
     for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){ 
      float dif_acum = As-B.elements[k*BLOCK_SIZE + featA]; 

      accum[featA] = dif_acum[featA]*dif_acum[featA]; 
      barrier(CLK_LOCAL_MEM_FENCE); 
      for(unsigned int i = 1; i < BLOCKSIZE; i *= 2) 
      { 
       if ((featA % (2*i)) == 0) 
       accum[featA] += accum[featA + i]; 
       barrier(CLK_LOCAL_MEM_FENCE); 
      } 
      if(featA == 0) 
       C[loop] = accum[0]; 
      barrier(CLK_LOCAL_MEM_FENCE); 

      loop += 1; 
     } 
    } 
} 

Answer 2

Grazie a Grizzly, che ho ora un kernel funzionante. Alcune cose che avevo bisogno di modificare in base alla risposta di Grizzly:

Ho aggiunto un'istruzione IF all'inizio della routine per scartare tutti i thread che non faranno riferimento ad alcuna posizione valida negli array che sto usando.

if(featA > BLOCK_SIZE){return;} 

Quando si copia il primo descrittore di memoria locale (comune) (I.G. a Bs), l'indice deve essere specificato poiché la funzione GetElement ritorna solo un elemento per ogni chiamata (ho saltato che la mia domanda).

Bs[featA] = GetElement(tmpA, 0, featA); 

Quindi, il ciclo SCAN necessario un piccolo ritocco perché il buffer viene sovrascritto dopo ogni iterazione e non si può controllare l'accesso quale thread i dati prima. Questo è il motivo per cui sto "riciclando" il buffer dif_acum per archiviare risultati parziali e in questo modo, prevenire incoerenze in tutto il ciclo.

dif_acum[featA] = accum[featA]; 

ci sono anche un certo controllo di confine nel ciclo SCAN per determinare in modo affidabile le condizioni per essere sommati.

if (featA >= j && next_addend >= 0 && next_addend < BLOCK_SIZE){ 

scorso, ho pensato che aveva un senso per includere l'incremento variabile di ciclo entro l'ultima IF in modo che solo un thread lo modifica.

if(featA == 0){ 
    C[loop] = accum[BLOCK_SIZE-1]; 
    loop += 1; 
} 

Questo è tutto. Mi chiedo ancora come posso utilizzare group_size per eliminare la definizione BLOCK_SIZE e se ci sono politiche migliori che posso adottare per quanto riguarda l'utilizzo dei thread.

Così il codice è, infine, come questo:

__kernel void CompareDescriptors(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float accum[BLOCK_SIZE], __local float Bs[BLOCK_SIZE]) 
{ 

    int gpidA = get_global_id(0); 
    int featA = get_local_id(0); 

    //global counter to store final differences 
    int loop = 0; 

    //auxiliary buffer to store temporary data 
    local float dif_acum[BLOCK_SIZE]; 

    //discard the threads that are not going to be used. 
    if(featA > BLOCK_SIZE){ 
     return; 
    } 

    //loop over all descriptors in A 
    for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){ 

     //take the gpidA-th descriptor 
     DescriptorList tmpA = GetDescriptor(A, i); 

     //copy the current descriptor to local memory 
     Bs[featA] = GetElement(tmpA, 0, featA); 

     //loop over all the descriptors in B 
     for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){ 
      //take the difference of both current descriptors 
      dif_acum[featA] = Bs[featA]-B.elements[k*BLOCK_SIZE + featA]; 

      //square the values in dif_acum 
      accum[featA] = dif_acum[featA]*dif_acum[featA]; 
      barrier(CLK_LOCAL_MEM_FENCE); 

      //copy the values of accum to keep consistency once the scan procedure starts. Mostly important for the first element. Two buffers are necesarry because the scan procedure would override values that are then further read if one buffer is being used instead. 
      dif_acum[featA] = accum[featA]; 

      //Compute the accumulated sum (a.k.a. scan) 
      for(int j = 1; j < BLOCK_SIZE; j *= 2){ 
       int next_addend = featA-(j/2); 
       if (featA >= j && next_addend >= 0 && next_addend < BLOCK_SIZE){ 
        dif_acum[featA] = accum[featA] + accum[next_addend]; 
       } 
       barrier(CLK_LOCAL_MEM_FENCE); 

       //copy As to accum 
       accum[featA] = GetElementArray(dif_acum, BLOCK_SIZE, featA); 
       barrier(CLK_LOCAL_MEM_FENCE); 
      } 

      //tell one of the threads to write the result of the scan in the array containing the results. 
      if(featA == 0){ 
       C[loop] = accum[BLOCK_SIZE-1]; 
       loop += 1; 
      } 
      barrier(CLK_LOCAL_MEM_FENCE); 

     } 
    } 
}