Berechnen von Partialsummen in OpenCL

Question

Ein 1D-Datensatz ist in Segmente unterteilt, jedes Arbeitselement verarbeitet ein Segment. Es liest eine Reihe von Elemente aus dem Segment? Die Anzahl der Elemente ist vorher nicht bekannt und unterscheidet sich für jedes Segment.

Zum Beispiel:

+----+----+----+----+----+----+----+----+----+  <-- segments 
    A BCD E FG HIJK L M  N  <-- elements in this segment 

Nachdem alle Segmente Prozesse gewesen sein sollten sie die elements in aneinander angrenzend Ausgabespeicher schreiben, wie

A B C D E F G H I J K L M N 

So ist die absolute Ausgangsposition der Elemente von einem Segment hängt von der Anzahl der Elemente in den vorherigen Segmenten ab. E befindet sich an Position 4, da Segment 1 Element (A) und Segment 2 3 Elemente enthält.

---

Der OpenCL-Kernel schreibt die Anzahl der Elemente für jedes Segment in einen lokalen/Shared-Memory-Puffer und funktioniert wie folgt (Pseudo-Code)

kernel void k(
    constant uchar* input, 
    global int* output, 
    local int* segment_element_counts 
) { 
    int segment = get_local_id(0); 
    int count = count_elements(&input[segment * segment_size]); 

    segment_element_counts[segment] = count; 

    barrier(CLK_LOCAL_MEM_FENCE); 

    ptrdiff_t position = 0; 
    for(int previous_segment = 0; previous_segment < segment; ++previous_segment) 
     position += segment_element_counts[previous_segment]; 

    global int* output_ptr = &output[position]; 
    read_elements(&input[segment * segment_size], output_ptr); 
} 

So wird jede Arbeitseinheit hat eine Teilsumme zu berechnen Verwenden einer Schleife, bei der die Arbeitselemente mit der größeren ID weitere Iterationen ausführen.

Gibt es eine effizientere Möglichkeit, dies zu implementieren (jedes Arbeitselement berechnet eine Teilsumme einer Sequenz bis zu seinem Index) in OpenCL 1.2? OpenCL 2 scheint dafür work_group_scan_inclusive_add bereitzustellen.

Answer 1

Sie tun können, teilweise N (Präfix) Summen in log2 (N) Iterationen mit so etwas wie folgt aus:

offsets[get_local_id(0)] = count; 
barrier(CLK_LOCAL_MEM_FENCE); 

for (ushort combine = 1; combine < total_num_segments; combine *= 2) 
{ 
    if (get_local_id(0) & combine) 
    { 
     offsets[get_local_id(0)] += 
      offsets[(get_local_id(0) & ~(combine * 2u - 1u)) | (combine - 1u)]; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 
} 

Da Segmentelement Grafen von

a  b  c  d 

Die aufeinanderfolgenden Iterationen produzieren:

a  b+a c  d+c 

und

a  b+a c+(b+a) (d+c)+(b+a) 

Welches ist das Ergebnis, das wir wollen.

In der ersten Iteration haben wir die Segmentelement-Zählungen in Gruppen von 2 aufgeteilt und summieren sie in ihnen. Dann führen wir 2 Gruppen gleichzeitig in 4 Elemente zusammen und propagieren das Ergebnis von der ersten Gruppe in die zweite. Wir wachsen die Gruppen wieder auf 8 und so weiter.

Die Schlüsselbeobachtung ist, dass dieses Muster passt auch die Binärdarstellung des Index jedes Segment:

0: 0b00 1: 0b01 2: 0b10 3: 0b11 

Index 0 führt keine Summen. Beide Indizes 1 und 3 führen in der ersten Iteration eine Summe durch (Bit 0/LSB = 1), während die Indizes 2 und 3 in der zweiten Iteration eine Summe durchführen (Bit 1 = 1).Das erklärt die folgende Zeile:

if (get_local_id(0) & combine) 

Die andere Erklärung, die wirklich eine Erklärung braucht, ist natürlich

 offsets[get_local_id(0)] += 
      offsets[(get_local_id(0) & ~(combine * 2u - 1u)) | (combine - 1u)]; 

Berechnung des Index, an dem wir die vorherigen Präfixsumme wir auf unsere Arbeit-Elements akkumulieren wollen finden Die Summe ist ein wenig schwierig. Der subexpression (combine * 2u - 1u) nimmt den Wert (2 n -1) bei jeder Iteration (für n beginnend bei 1):

1 = 0b001 
3 = 0b011 
7 = 0b111 
… 

Durch bitweise maskierende Bit diese Suffixe aus (dh i & ~x) das Werkelementindex, Dies gibt Ihnen den Index der ersten Element in der aktuellen Gruppe.

Der Unterausdruck (combine - 1u) gibt Ihnen dann den Index innerhalb der aktuellen Gruppe der letzten Artikel der ersten Hälfte. Wenn Sie die beiden zusammenfügen, erhalten Sie den Gesamtindex des Artikels, den Sie im aktuellen Segment ansammeln möchten.

Es gibt eine leichte Hässlichkeit im Ergebnis: Es ist um eins nach links verschoben: so muss Segment 1 offsets[0] usw. verwenden, während der Offset von Segment 0 natürlich 0 ist. Sie können die Offsets entweder überzuordnen Array um 1 und führen Sie die Präfixsummen auf dem Subarray beginnend bei Index 1 aus und initialisieren Sie den Index 0 auf 0 oder verwenden Sie eine Bedingung.

Es gibt wahrscheinlich Profiling-gesteuerte Mikro-Optimierungen, die Sie an dem obigen Code vornehmen können.