2010-09-22 2 views
5

Ich berechne den Euklidischen Abstand zwischen n-dimensionalen Punkten mit OpenCL. Ich bekomme zwei Listen von n-dimensionalen Punkten und ich sollte ein Array zurückgeben, das nur die Abstände von jedem Punkt in der ersten Tabelle zu jedem Punkt in der zweiten Tabelle enthält.Kumulative Array-Summierung mit OpenCL

Mein Ansatz ist die regelmäßige doble Schleife zu tun (für jeden Punkt in Tabelle 1 {für jeden Punkt in Tabelle 2 {...}} und führen Sie dann die Berechnung für jedes Paar von Punkten in paralell.

Die euklidischen Der Abstand wird dann in 3 Teile geteilt: 1. nehmen Sie den Unterschied zwischen den einzelnen Punkten in den Punkten 2. Quadrat dieser Unterschied (immer noch für jede Dimension) 3. addieren Sie alle Werte in 2. 4. Nehmen Sie die Quadratwurzel des Wertes, der in 3 erhalten wurde. (Dieser Schritt wurde in diesem Beispiel weggelassen.)

Alles funktioniert wie ein Zauber bis ich versuche, die Summe aller Differenzen zu akkumulieren (d. h. Ausführung von Schritt 3. des oben beschriebenen Verfahrens, Zeile 49 des folgenden Codes).

Als Testdaten verwende ich DescriptorLists mit je 2 Punkten: DescriptorList1: 001,002,003, ..., 127,128; (p1) 129,130,131, ..., 255,256; (p2)

Deskriptorliste2: 000,001,002, ..., 126,127; (p1) 128,129,130, ..., 254,255; (P2)

So ist der resultierende Vektor sollte die Werte: 128, 2064512, 2130048, 128 Im Moment bin ich immer Zufallszahlen, die mit jedem Lauf variieren.

Ich freue mich über jede Hilfe oder führt, was ich falsch mache. Hoffentlich ist alles klar über das Szenario in arbeite ich

#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); 
     } 
    } 
} 

Antwort

7

Ihr Problem in diesen Zeilen Code liegt.

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

Alle Themen in Ihrer Arbeitsgruppe (na ja, eigentlich alle Threads, aber lasst uns später dazu kommen) versuchen, diese Array-Position gleichzeitig ohne jegliche Synchronisation zu modifizieren. Mehrere Faktoren machen das wirklich problematisch:

  1. Die Arbeitsgruppe garantiert nicht vollständig parallel zu arbeiten, dass für einige Threads Bedeutung C [loop] = 0 nach anderen Threads bereits die nächste Zeile
  2. ausgeführt haben aufgerufen werden können
  3. Diejenigen, die parallel ausgeführt werden, lesen den gleichen Wert von C [loop], modifizieren sie mit ihrem Inkrement und versuchen, auf dieselbe Adresse zurück zu schreiben. Ich bin mir nicht ganz sicher, was das Ergebnis dieses Writebacks ist (ich denke, dass einer der Threads es schafft, zurückzuschreiben, während die anderen scheitern, aber ich bin mir nicht ganz sicher), aber es ist in jedem Fall falsch.

Jetzt können dieses Problem beheben: Während wir diese bekommen vielleicht in der Lage auf den globalen Speicher zu arbeiten atomics verwenden, wird es nicht schnell sein, können im lokalen Speicher so akkumulieren:

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]; 

Of Natürlich können Sie andere lokale Puffer dafür verwenden, aber ich denke, der Punkt ist klar (übrigens: Sind Sie sicher, dass dif_acum im lokalen Speicher erstellt wird, weil ich denke, ich habe irgendwo gelesen, dass dies nicht im lokalen Speicher abgelegt würde, was würde das Vorladen von A in den lokalen Speicher sinnlos machen).

Einige andere Punkte zu diesem Code:

  1. Ihr Code scheint auf die Verwendung nur auf Arbeitsgruppe ausgerichtet zu werden (Sie verwenden entweder groupid noch global-ID nicht, welche Elemente zu sehen arbeiten), für optimale Leistung möchten Sie vielleicht mehr als das verwenden.
  2. Könnte persönliche preferance sein, aber ich mir scheint es besser get_local_size(0) für die workgroupsize zu benutzen als ein Definieren zu verwenden
  3. (da Sie es in dem Host-Code ändern, könnten Sie ohne zu merken, Ihren OpenCL-Code geändert haben sollte) Die Barrieren in Ihrem Code sind alles unnötig, da kein Thread auf ein Element im lokalen Speicher zugreift, das von einem anderen Thread geschrieben wird. Daher müssen Sie hierfür keinen lokalen Speicher verwenden.

die letzte Kugel Betrachtet man einfach tun konnte:

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

Dies würde den Code machen (nicht die ersten beiden Kugeln bedenkt):

__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; 
     } 
    } 
} 
+0

ich zu sagen habe, vor allem, dass ich mit Antwort des Grizzly sehr dankbar bin. Ich bin ziemlich neu in OpenCL, und obwohl ich den Beispielcode, den er ein bisschen gegeben hat, anpassen musste, führte es mich direkt in die richtige Richtung.Wichtige Dinge, die mir aufgefallen sind (durch Versuch und Irrtum): Threads, die die Array-Positionen nicht adressieren, müssen verworfen werden; Die SCAN-Schleife erforderte ein wenig Feinabstimmung, nämlich die Verwendung eines Hilfspuffers, um Teilergebnisse zu akkumulieren und nach Randbedingungen zu suchen, um die hinzuzufügenden Terme zu finden. Danke dir nochmal! Ich poste den Code, der für mich funktioniert hat. – SebastianP

3

Dank Grizzly, habe ich jetzt ein funktionierender Kernel. Einige Dinge, die ich in der Antwort von Grizzly ändern musste:

Ich habe eine IF-Anweisung am Anfang der Routine hinzugefügt, um alle Threads zu verwerfen, die keine gültige Position in den Arrays, die ich verwende, referenzieren.

if(featA > BLOCK_SIZE){return;} 

Beim Kopieren des ersten Deskriptors auf lokale (gemeinsam) Speicher (i.g. zu Bs), der Index angegeben werden muss, da die Funktion GetElement pro Anruf nur ein Element gibt (ich übersprungen, dass auf meiner Frage).

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

Dann benötigt die Abtastschleife ein wenig Feintuning, da der Puffer nach jeder Iteration überschrieben wird, und man kann die Daten zunächst keinen Zugriff steuern, welcher Thread. Deshalb "recycle" ich den dif_acum-Puffer, um Teilergebnisse zu speichern und auf diese Weise Inkonsistenzen in dieser Schleife zu verhindern.

Es gibt auch einige Grenzkontrolle in der SCAN-Schleife, um die zu addierenden Terme zuverlässig zu bestimmen.

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

Last, ich dachte, es sinnvoll, den Schleifenvariable Schritt innerhalb der letzten aufzunehmen gemacht IF-Anweisung, so dass nur ein Thread es modifiziert.

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

Das ist es. Ich frage mich immer noch, wie kann ich group_size verwenden, um diese BLOCK_SIZE-Definition zu eliminieren, und wenn es bessere Richtlinien gibt, die ich bezüglich der Thread-Verwendung annehmen kann.

So sieht der Code schließlich wie folgt aus:

__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); 

     } 
    } 
} 
+0

Ich denke immer noch, dass Sie diese lokalen Puffer nicht brauchen (accept for accum natürlich), da sowohl dif_acum als auch Bs nur auf Position featA zugegriffen werden, die die lokale ID des Threads ist und daher auf jedes Element der Arrays zugegriffen wird nur ein Thread. Außerdem sollte die Verwendung von zwei Puffern für die Scan-Schleife nicht wirklich notwendig sein, da die Konsistenz durch die Barrieren gewährleistet ist (die Iterationen sind durch Barrieren getrennt und werden bei jeder Iteration nur auf die Elemente zugegriffen, auf die jeweils zugegriffen wird) ein Thread – Grizzly