Hybridimplementierung des MST-Algorithmus mit CPU und GPU

    Einleitung


    Die Lösung des Problems der Suche nach minimalen Spannbäumen (MST - Minimum Spanning Tree) ist eine häufige Aufgabe in verschiedenen Forschungsbereichen: Erkennung verschiedener Objekte, Computer Vision, Analyse und Aufbau von Netzwerken (z. B. Telefon, Elektrik, Computer, Straße usw.), Chemie und Biologie und viele andere. Es gibt mindestens drei bekannte Algorithmen, die dieses Problem lösen: Boruvki, Kruskal und Prima. Das Verarbeiten von großen Grafiken (die mehrere GB belegen) ist für die Zentraleinheit (CPU) eine recht zeitaufwendige Aufgabe, die derzeit gefragt ist. Grafikbeschleuniger (Graphic Accelerators, GPUs), die eine wesentlich höhere Leistung als CPUs aufweisen können, finden zunehmend Verbreitung. Aber die MST-Task passt, wie viele Grafikverarbeitungsaufgaben, nicht gut zur GPU-Architektur. Dieser Artikel beschreibt die Implementierung dieses Algorithmus auf der GPU. Es wird auch gezeigt, wie die CPU verwendet werden kann, um eine hybride Implementierung dieses Algorithmus auf dem gemeinsam genutzten Speicher eines Knotens (bestehend aus einer GPU und mehreren CPUs) zu erstellen.


    Beschreibung des Grafikformats


    Wir betrachten kurz die Speicherstruktur eines ungerichteten gewichteten Graphen, da dieser in Zukunft erwähnt und transformiert wird. Das Diagramm wird in komprimierter CSR (Compressed Sparse Row) angegeben. [1]formatieren. Dieses Format wird häufig zum Speichern von dünnen Matrizen und Diagrammen verwendet. Für ein Diagramm mit N Eckpunkten und M Kanten werden drei Arrays benötigt: X, A und W. Bei einem Array X der Größe N + 1 sind die beiden anderen 2 * M, da in einem nicht orientierten Diagramm für jedes Paar von Eckpunkten die direkten und umgekehrten Bögen gespeichert werden müssen. Array X speichert den Anfang und das Ende der Liste der Nachbarn, die in Array A gespeichert sind, dh die gesamte Liste der Nachbarn des Scheitelpunkts J befindet sich in Array A von Index X [J] bis X [J + 1], ohne diesen einzuschließen. Die ähnlichen Indizes speichern die Gewichte jeder Kante vom Scheitelpunkt J. Zur Veranschaulichung zeigt die folgende Abbildung das Diagramm mit 6 unter Verwendung der Adjazenzmatrix geschriebenen Scheitelpunkten und dem CSR-Format auf der rechten Seite (der Einfachheit halber ist das Gewicht jeder Kante nicht angegeben).


    Getestete Grafiken


    Ich werde sofort beschreiben, auf welchen Graphen das Testen stattgefunden hat, da eine Beschreibung der Transformationsalgorithmen und des MST-Algorithmus Kenntnisse über die Struktur der fraglichen Graphen erfordert. Zur Bewertung der Implementierungsleistung werden zwei Arten von synthetischen Diagrammen verwendet: RMAT-Diagramme und SSCA2-Diagramme. R-MAT-Graphen modellieren gut reale Graphen aus sozialen Netzwerken, dem Internet [2] . In diesem Fall betrachten wir RMAT-Graphen mit einem durchschnittlichen Verknüpfungsgrad des Scheitelpunkts 32, und die Anzahl der Scheitelpunkte ist eine Zweierpotenz. In einem solchen RMAT-Graphen gibt es eine große zusammenhängende Komponente und eine Anzahl kleiner zusammenhängender Komponenten oder hängender Eckpunkte. SSCA2-Graph ist eine große Menge unabhängiger Komponenten, die durch Kanten miteinander verbunden sind. [3]. Ein SSCA2-Diagramm wird generiert, sodass der durchschnittliche Konnektivitätsgrad des Scheitelpunkts nahe bei 32 liegt und die Anzahl der Scheitelpunkte ebenfalls eine Zweierpotenz ist. Somit werden zwei völlig unterschiedliche Graphen betrachtet.

    Eingabekonvertierung


    Da der Test des Algorithmus an den Graphen RMAT und SSCA2 durchgeführt wird, die unter Verwendung des Generators erhalten werden, sind einige Transformationen erforderlich, um die Leistung des Algorithmus zu verbessern. Alle Conversions werden nicht in die Leistungsberechnung einbezogen.

    1. Lokale Sortierung der Scheitelpunktliste
      Für jeden Scheitelpunkt sortieren wir die Liste der Nachbarn nach Gewicht in aufsteigender Reihenfolge. Dies vereinfacht teilweise die Wahl der Mindestkante bei jeder Iteration des Algorithmus. Da diese Sortierung lokal ist, bietet sie keine vollständige Lösung für das Problem.
    2. Alle Scheitelpunkte des Diagramms neu
      nummerieren Wir nummerieren die Scheitelpunkte des Diagramms so, dass die am meisten verbundenen Scheitelpunkte die nächstgelegenen Nummern haben. Infolge dieser Operation ist in jeder verbundenen Komponente die Differenz zwischen der maximalen und der minimalen Scheitelpunktzahl die kleinste, wodurch der Cache für kleine Grafikprozesse optimal genutzt werden kann. Es ist anzumerken, dass diese Umnummerierung für RMAT-Diagramme keinen signifikanten Effekt hat, da dieses Diagramm eine sehr große Komponente hat, die auch nach Anwendung dieser Optimierung nicht in den Cache passt. Bei SSCA2-Diagrammen ist der Effekt dieser Transformation deutlicher, da in diesem Diagramm eine große Anzahl kleiner Komponenten vorhanden ist.
    3. Anzeigen eines Graphen Gewichte auf ganze Zahlen
      In diesem Problem brauchen wir keine Transaktionen in der Bilanz Diagramm zu machen. Wir müssen in der Lage sein, die Gewichte von zwei Rippen zu vergleichen. Für diese Zwecke können Sie Ganzzahlen anstelle von Zahlen mit doppelter Genauigkeit verwenden, da die Verarbeitungsgeschwindigkeit von Zahlen mit einfacher Genauigkeit auf der GPU viel höher ist als die von Zahlen mit doppelter Genauigkeit. Diese Transformation kann für Diagramme durchgeführt werden, deren Anzahl eindeutiger Kanten 2 ^ 32 nicht überschreitet (die maximale Anzahl unterschiedlicher Zahlen, die in ein vorzeichenloses int passen). Wenn der durchschnittliche Konnektivitätsgrad jedes Scheitelpunkts 32 m beträgt, hat der größte Graph, der mit dieser Transformation verarbeitet werden kann, 2 ^ 28 Scheitelpunkte und belegt 64 GB Speicher. Heute ist die größte Menge an Speicher in Tesla K40 Beschleunigern NVidia [4] / NVidia Titan X[5] und AMD FirePro w9100 [6] haben 12 GB bzw. 16 GB. Daher können Sie auf einer einzelnen GPU mit dieser Konvertierung sehr große Diagramme verarbeiten.
    4. Vertex-Komprimierung
      Diese Transformation ist aufgrund ihrer Struktur nur auf SSCA2-Diagramme anwendbar. Bei dieser Aufgabe spielt die Speicherleistung aller Ebenen eine entscheidende Rolle: vom globalen Speicher bis zum Cache der ersten Ebene. Um den Datenverkehr zwischen dem globalen Speicher und dem L2-Cache zu reduzieren, können Sie Scheitelpunktinformationen in komprimierter Form speichern. Anfänglich werden Informationen zu Scheitelpunkten in Form von zwei Arrays dargestellt: Array X, in dem der Anfang und das Ende der Liste der Nachbarn in Array A gespeichert sind (ein Beispiel für nur einen Scheitelpunkt):

      Scheitelpunkt J hat 10 Nachbarscheitelpunkte, und wenn die Nummer jedes Nachbarn unter Verwendung des vorzeichenlosen int-Typs gespeichert wird, sind 10 * sizeof (vorzeichenlose int) Bytes erforderlich, um die Liste der Nachbarn des J-Scheitelpunkts zu speichern, und 2 * M * sizeof (für das gesamte Diagramm vorzeichenlos). int) byte. Wir nehmen an, dass sizeof (unsigned int) = 4 Byte, sizeof (unsigned short) = 2 Byte, sizeof (unsigned char) = 1 Byte ist. Dann werden für diesen Eckpunkt 40 Bytes benötigt, um die Liste der Nachbarn zu speichern.
      Es ist nicht schwer zu bemerken, dass der Unterschied zwischen der maximalen und der minimalen Anzahl der Eckpunkte in dieser Liste 8 beträgt und nur 4 Bits benötigt werden, um diese Anzahl zu speichern. Basierend auf diesen Überlegungen, dass der Unterschied zwischen der maximalen und der minimalen Anzahl eines Scheitelpunkts kleiner als das vorzeichenlose int sein kann, können wir die Anzahl jedes Scheitelpunkts wie folgt darstellen:
      base_J + 256 * k + short_endV,
      Dabei ist base_J beispielsweise die minimale Scheitelpunktzahl aus der gesamten Liste der Nachbarn. In diesem Beispiel ist es 1. Diese Variable hat den Typ int ohne Vorzeichen und es gibt so viele solche Variablen, wie Scheitelpunkte im Diagramm vorhanden sind. Als nächstes berechnen wir die Differenz zwischen der Scheitelpunktzahl und der ausgewählten Basis. Da wir den kleinsten Peak als Basis gewählt haben, ist dieser Unterschied immer positiv. Für SSCA2 wird dieser Unterschied in ein Short ohne Vorzeichen gesetzt. short_endV ist der Rest der Division durch 256. Zum Speichern dieser Variablen verwenden wir den vorzeichenlosen char-Typ. und k ist der ganzzahlige TeilWählen Sie für k 2 Bits aus (dh k liegt im Bereich von 0 bis 3). Die gewählte Darstellung reicht für die betreffende Grafik aus. In der bitweisen Darstellung sieht es so aus: Um die

      Liste der Eckpunkte zu speichern, ist für dieses Beispiel (1 + 0,25) * 10 + 4 = 16,5 Bytes anstelle von 40 Bytes und für den gesamten Graphen erforderlich: (2 * M + 4 * N + 2 * M / 4) anstelle von 2 * M * 4. Wenn N = 2 * M / 32, verringert sich die Gesamtlautstärke in
      (8 * M) / (2 * M + 8 * M / 32 + 2 * M / 4) = 2,9 mal


    Allgemeine Beschreibung des Algorithmus


    Um den MST-Algorithmus zu implementieren, wurde der Boruwka-Algorithmus gewählt. Die grundlegende Beschreibung des Boruwka-Algorithmus und eine Darstellung seiner Iterationen sind hier gut dargestellt [7] .
    Gemäß dem Algorithmus sind anfangs alle Scheitelpunkte im Minimalbaum enthalten. Als Nächstes müssen Sie die folgenden Schritte ausführen:
    1. Finden Sie die minimalen Kanten zwischen allen Bäumen für die nachfolgende Vereinigung. Wenn in diesem Schritt keine Kanten ausgewählt sind, wird die Antwort auf das Problem empfangen
    2. Füge passende Bäume zusammen. Dieser Schritt ist in zwei Phasen unterteilt: Löschen von Zyklen, da sich zwei Bäume gegenseitig als Kandidaten für das Zusammenführen angeben können, und der Zusammenführungsschritt, wenn die Nummer des Baums ausgewählt wird, der die kombinierten Teilbäume enthält. Zur Sicherheit wählen wir die Mindestanzahl. Wenn während der Zusammenführung nur noch ein Baum übrig bleibt, wird die Antwort auf das Problem empfangen.
    3. Nummerieren Sie die resultierenden Bäume neu, um mit dem ersten Schritt fortzufahren (sodass alle Bäume Zahlen von 0 bis k haben).


    Algorithmusstufen


    Im Allgemeinen lautet der implementierte Algorithmus wie folgt:

    Der Austritt aus dem gesamten Algorithmus erfolgt in zwei Fällen: Wenn alle Eckpunkte nach N Iterationen zu einem Baum zusammengefasst werden oder wenn es unmöglich ist, die Mindestkante von jedem Baum zu finden (in diesem Fall werden die Mindestüberspannungsbäume gefunden).

    1. Suchen Sie nach der Mindestkante.


    Zunächst wird jeder Scheitelpunkt des Diagramms in einem separaten Baum platziert. Als nächstes findet ein iterativer Prozess zum Kombinieren von Bäumen statt, der aus den vier oben diskutierten Prozeduren besteht. Mit dem Verfahren zum Ermitteln der minimalen Kante können Sie genau die Kanten auswählen, die im minimalen Spannbaum enthalten sein sollen. Wie oben beschrieben, wird bei der Eingabe dieser Prozedur der konvertierte Graph im CSR-Format gespeichert. Da die Kanten für die Liste der Nachbarn teilweise nach Gewicht sortiert waren, wird durch Auswahl des minimalen Scheitelpunkts nur die Liste der Nachbarn angezeigt und der erste Scheitelpunkt ausgewählt, der zu einem anderen Baum gehört. Wenn wir annehmen, dass das Diagramm keine Schleifen enthält, reduziert sich die Auswahl des minimalen Scheitelpunkts im ersten Schritt des Algorithmus auf die Auswahl des ersten Scheitelpunkts aus der Liste der Nachbarn für jeden betrachteten Scheitelpunkt. weil die Liste benachbarter Scheitelpunkte (die zusammen mit dem betreffenden Scheitelpunkt die Kanten des Diagramms bilden) durch Erhöhen des Gewichts der Kanten sortiert wird und jeder Scheitelpunkt in einem separaten Baum enthalten ist. In jedem anderen Schritt müssen Sie die Liste aller benachbarten Scheitelpunkte der Reihe nach anzeigen und den Scheitelpunkt auswählen, der zu einem anderen Baum gehört.

    Warum ist es unmöglich, den zweiten Scheitelpunkt aus der Liste der benachbarten Scheitelpunkte auszuwählen und diese Kante minimal zu setzen? Nach der Prozedur des Kombinierens von Bäumen (die später betrachtet wird) kann es vorkommen, dass einige Scheitelpunkte aus der Liste benachbarter Scheitelpunkte im selben Baum wie der betrachtete erscheinen. Dadurch wird diese Kante zu einer Schleife für diesen Baum, und gemäß den Bedingungen des Algorithmus muss die Mindestkante ausgewählt werden zu anderen Bäumen.

    Union Find eignet sich gut zum Implementieren der Vertex-Verarbeitung und zum Durchführen der Listen für Suchen, Zusammenführen und Zusammenführen. [8]. Leider werden nicht alle Strukturen auf der GPU optimal verarbeitet. Am rentabelsten bei dieser Aufgabe ist (wie bei den meisten anderen) die Verwendung von fortlaufenden Arrays im GPU-Speicher anstelle von verknüpften Listen. Im Folgenden werden ähnliche Algorithmen zum Ermitteln der Mindestkante, zum Kombinieren von Segmenten und zum Löschen von Zyklen im Diagramm betrachtet.

    Betrachten Sie den Algorithmus zum Ermitteln der Mindestkante. Es kann in zwei Schritten dargestellt werden:
    • Auswahl der Mindestkante, die von jedem Scheitelpunkt (der in einem bestimmten Segment enthalten ist) des betreffenden Graphen ausgeht;
    • Auswahl von Kanten mit minimalem Gewicht für jeden Baum.

    Um die im CSR-Format aufgezeichneten Scheitelpunktinformationen nicht zu verschieben, werden zwei Hilfsarrays verwendet, die den Index des Anfangs und des Endes von Array A der Nachbarliste speichern. Zwei gegebene Arrays geben Segmente von Vertex-Listen an, die zu einem Baum gehören. Zum Beispiel wird im ersten Schritt das Array von Anfängen oder niedrigeren Werten die Werte 0..N des Arrays X und das Array von Enden oder die oberen Werte die Werte 1..N + 1 des Arrays X haben. Und dann nach der Prozedur des Kombinierens von Bäumen (die später betrachtet wird ) werden diese Segmente gemischt, aber das Array der Nachbarn A wird im Speicher nicht geändert.

    Beide Schritte können parallel ausgeführt werden. Um den ersten Schritt abzuschließen, müssen Sie sich die Liste der Nachbarn jedes Scheitelpunkts (oder jedes Segments) ansehen und die erste Kante auswählen, die zu einem anderen Baum gehört. Sie können einen Warp (bestehend aus 32 Threads) auswählen, um die Liste der Nachbarn für jeden Vertex anzuzeigen. Es sei daran erinnert, dass mehrere Segmente des Arrays benachbarter Peaks A nicht in einer Reihe liegen können und zu einem Baum gehören (die Segmente von Baum 0 sind rot und der Baum 1 grün hervorgehoben):

    Da jedes Segment der Liste der Nachbarn sortiert ist, müssen nicht alle Eckpunkte angezeigt werden. Da ein Warp aus 32 Threads besteht, erfolgt die Anzeige in Teilen von 32 Vertices. Nachdem 32 Scheitelpunkte angezeigt wurden, muss das Ergebnis kombiniert werden. Wenn nichts gefunden wird, werden die nächsten 32 Scheitelpunkte angezeigt. Um das Ergebnis zu kombinieren, können Sie den Scan-Algorithmus verwenden [9] . Es ist möglich, diesen Algorithmus in einem Warp zu implementieren, indem Shared Memory oder neue shfl-Anweisungen verwendet werden. [10](verfügbar von der Kepler-Architektur), die den Datenaustausch zwischen den Threads eines Warps für einen Befehl ermöglicht. Als Ergebnis der Experimente stellte sich heraus, dass shfl-Anweisungen die Arbeit des gesamten Algorithmus um etwa das Zweifache beschleunigen können. Somit kann diese Operation unter Verwendung von shfl-Anweisungen ausgeführt werden, zum Beispiel wie folgt:
    	unsigned idx = blockIdx.x * blockDim.x + threadIdx.x; // глобальный индекс нити
    	unsigned lidx = idx % 32;  
    #pragma unroll
    	for (int offset = 1; offset <= 16; offset *= 2)
    	{
    		tmpv = __shfl_up(val, (unsigned)offset);
    		if(lidx >= offset)
    			val += tmpv;
    	}
    	tmpv = __shfl(val, 31); // рассылка всем нитям последнего значения. Если получено значение 1, то какая-то нить нашла 
                                // минимальное ребро, иначе необходимо продолжить поиск.
    

    Als Ergebnis dieses Schritts werden die folgenden Informationen für jedes Segment aufgezeichnet: Die Nummer des Scheitelpunkts in Array A, der in der Kante mit dem Mindestgewicht enthalten ist, und die Gewicht der Kante selbst. Wenn nichts gefunden wird, kann zum Beispiel die Zahl N + 2 in die Scheitelpunktzahl geschrieben werden.

    Der zweite Schritt ist notwendig, um die ausgewählten Informationen zu reduzieren, nämlich die Auswahl der Kanten mit einem Mindestgewicht für jeden der Bäume. Dieser Schritt wird ausgeführt, weil Segmente, die zu demselben Baum gehören, parallel und unabhängig voneinander gescannt werden und für jedes Segment eine minimale Gewichtungskante ausgewählt wird. In diesem Schritt kann ein Warp die Informationen für jeden Baum (für mehrere Segmente) reduzieren, und shfl-Anweisungen können auch zur Reduzierung verwendet werden. Nach Abschluss dieses Schritts ist bekannt, mit welchem ​​Baum jeder der Bäume durch eine minimale Kante verbunden ist (falls vorhanden). Um diese Informationen aufzuzeichnen, führen wir zwei weitere Hilfsarrays ein, von denen eines die Anzahl der Bäume speichert, bis zu denen es eine Mindestkante gibt, und das zweite die Anzahl der Scheitelpunkte im Originaldiagramm, die die Wurzel der Scheitelpunkte sind, die in den Baum eintreten.

    Es ist zu beachten, dass für die Arbeit mit Indizes zwei weitere Arrays erforderlich sind, mit deren Hilfe die ursprünglichen Indizes in neue Indizes konvertiert werden und der ursprüngliche Index mithilfe des neuen Index abgerufen werden kann. Diese sogenannten Indexkonvertierungstabellen werden bei jeder Iteration des Algorithmus aktualisiert. Die Tabelle zum Erhalten eines neuen Index durch den Anfangsindex hat die Größe N - die Anzahl der Eckpunkte in der Grafik, und die Tabelle zum Erhalten des neuen Index wird mit jeder Iteration auf eine neue Weise reduziert und hat eine Größe, die gleich der Anzahl der Bäume in jeder gewählten Iteration des Algorithmus ist (bei der ersten Iteration des Algorithmus hat diese Tabelle auch Größe N).

    2. Zyklen löschen.


    Dieses Verfahren ist erforderlich, um Schleifen zwischen zwei Bäumen zu entfernen. Diese Situation tritt auf, wenn der Baum N1 eine minimale Kante zum Baum N2 hat und der Baum N2 die minimale Kante zum Baum N1 hat. In der obigen Abbildung gibt es nur einen Zyklus zwischen zwei Bäumen mit den Nummern 2 und 4. Da bei jeder Iteration weniger Bäume vorhanden sind, wählen wir die Mindestanzahl von zwei Bäumen, aus denen der Zyklus besteht. In diesem Fall zeigt 2 auf 2 und 4 zeigt weiterhin auf 2. Mit diesen Überprüfungen können Sie einen solchen Zyklus bestimmen und zugunsten der Mindestanzahl beseitigen:
        unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
        unsigned local_f = сF[i]; 
        if (сF[local_f] == i)
        {
             if (i < local_f)
             {
                    F[i] = i;
                    . . . . . . .  
              }
        }
    

    Diese Prozedur kann parallel ausgeführt werden, da jeder Scheitelpunkt unabhängig verarbeitet werden kann und sich Datensätze in einem neuen Array von Scheitelpunkten ohne Zyklen nicht schneiden.

    3. Die Vereinigung der Bäume.


    Diese Prozedur kombiniert die Bäume zu größeren. Die Prozedur zum Löschen von Schleifen zwischen zwei Bäumen ist im Wesentlichen eine Vorverarbeitung vor dieser Prozedur. Es vermeidet Schleifen beim Zusammenführen von Bäumen. Beim Kombinieren von Bäumen wird eine neue Wurzel ausgewählt, indem Verknüpfungen geändert werden. Wenn beispielsweise Baum 0 auf Baum 1 und Baum 1 wiederum auf Baum 3 zeigt, können Sie die Verknüpfung von Baum 0 von Baum 1 auf Baum 3 ändern. Diese Verknüpfungsänderung ist sinnvoll, wenn die Verknüpfung nicht zu einer Schleife zwischen den beiden Bäumen führt . In Anbetracht des obigen Beispiels bleibt nach dem Löschen der Zyklen und der Vereinigung der Bäume nur ein Baum mit der Nummer 2. Der Vereinigungsprozess kann folgendermaßen dargestellt werden:

    Die Struktur des Graphen und das Prinzip seiner Verarbeitung ist so, dass es keine Situation gibt, in der die Prozedur wiederholt wird, und dass sie auch parallel ausgeführt werden kann.

    4. Peaks (Bäume) neu nummerieren.


    Nach dem Zusammenführen müssen die resultierenden Bäume neu nummeriert werden, sodass ihre Nummern in einer Reihe von 0 bis P liegen. Konstruktionsbedingt müssen die neuen Nummern Array-Elemente erhalten, die die Bedingung F [i] == i erfüllen (für das oben betrachtete Beispiel erfüllt nur das Element diese Bedingung mit Index 2). Mit atomaren Operationen können Sie also das gesamte Array mit neuen Werten von 1 ... (P + 1) markieren. Vervollständigen Sie als Nächstes die Tabellen, um einen neuen Index mit dem Anfangsindex und dem Anfangsindex zu erhalten. Das

    Arbeiten mit diesen Tabellen wird in der Prozedur zum Ermitteln der Mindestkante beschrieben. Die nächste Iteration kann nicht korrekt ausgeführt werden, ohne die Tabellendaten zu aktualisieren. Alle beschriebenen Vorgänge werden parallel und auf der GPU ausgeführt.

    Um es zusammenzufassen. Alle 4 Vorgänge werden parallel und auf einem Grafikbeschleuniger durchgeführt. Es wird mit eindimensionalen Arrays gearbeitet. Die einzige Schwierigkeit besteht darin, dass alle diese Prozeduren indirekt indiziert werden. Und um Cache-Ausfälle bei der Arbeit mit Arrays zu reduzieren, wurden verschiedene Permutationen des eingangs beschriebenen Graphen verwendet. Leider gelingt es nicht jedem Graphen, Verluste durch indirekte Indizierung zu reduzieren. Wie später gezeigt wird, erzielen RMAT-Diagramme mit diesem Ansatz keine sehr hohe Leistung. Das Ermitteln der minimalen Flanke nimmt bis zu 80% der Zeit in Anspruch, während der Rest die verbleibenden 20% ausmacht. Dies liegt an der Tatsache, dass bei den Prozeduren zum Kombinieren, Löschen von Schleifen und Umnummerieren von Scheitelpunkten mit Arrays gearbeitet wird, deren Länge ständig abnimmt (von Iteration zu Iteration). Für die betrachteten Graphen müssen ca. 7-8 Iterationen durchgeführt werden. Dies bedeutet, dass die Anzahl der bereits im ersten Schritt zu verarbeitenden Eckpunkte viel geringer als N / 2 wird. Während im Hauptverfahren zum Auffinden der Mindestkante mit Anordnungen von Eckpunkten A und einer Anordnung von Gewichten W gearbeitet wird (obwohl bestimmte Elemente ausgewählt sind).
    Zusätzlich zum Speichern des Graphen wurden mehrere Arrays der Länge N verwendet:
    • ein Array von unteren Werten und ein Array von oberen Werten. Wird verwendet, um mit Segmenten von Array A zu arbeiten.
    • Array-Tabelle, um den ursprünglichen Index für einen neuen zu erhalten;
    • Array-Tabelle, um einen neuen Index für das Original zu erhalten;
    • ein Array für Scheitelpunktnummern und ein Array für die diesen entsprechenden Gewichte, die im zweiten Schritt des Minimum-Kanten-Suchvorgangs verwendet werden;
    • Hilfsarray, mit dem Sie im ersten Schritt des Verfahrens bestimmen können, zu welcher Mindestkante der Baum dieses oder jenes Segment gehört.


    Hybride Implementierung des Minimum Edge Search-Verfahrens.



    Der oben beschriebene Algorithmus arbeitet auf einer einzelnen GPU nicht schlecht. Die Lösung dieses Problems ist so organisiert, dass Sie versuchen können, diesen Vorgang auch auf der CPU zu parallelisieren. Dies ist natürlich nur mit Shared Memory möglich. Dazu habe ich den OpenMP-Standard verwendet und Daten zwischen CPU und GPU über den PCIe-Bus übertragen. Wenn wir uns die Ausführung von Prozeduren in einer Iteration auf der Zeitachse vorstellen, sieht das Bild bei Verwendung einer GPU ungefähr so ​​aus:

    Zunächst werden alle Grafikdaten sowohl in der CPU als auch in der GPU gespeichert. Damit die CPU lesen kann, müssen Informationen über Segmente übertragen werden, die beim Zusammenführen von Bäumen verschoben wurden. Damit die GPU die Iteration des Algorithmus fortsetzt, müssen die berechneten Daten zurückgegeben werden. Es wäre logisch, asynchrones Kopieren zwischen dem Host und dem Beschleuniger zu verwenden:

    Der Algorithmus auf der CPU wiederholt den auf der GPU verwendeten Algorithmus, nur OpenMP wird zur Parallelisierung der Schleife verwendet. [11]. Wie erwartet zählt die CPU nicht so schnell wie die GPU, und der Kopieraufwand stört ebenfalls. Damit die CPU ihren Teil berechnen kann, müssen die Berechnungsdaten im Verhältnis 1: 5 geteilt werden, dh es sollten nicht mehr als 20% -25% an die CPU übertragen werden, und der Rest sollte auf der GPU berechnet werden. Der Rest der Prozeduren ist hier und da nicht vorteilhaft zu lesen, da sie sehr wenig Zeit in Anspruch nehmen und der Overhead und die langsame CPU-Geschwindigkeit nur die Algorithmuszeit erhöhen. Die Kopiergeschwindigkeit zwischen CPU und GPU ist ebenfalls sehr wichtig. Die getestete Plattform unterstützte PCIe 3.0, wodurch 12 GB / s erreicht wurden.

    Bisher unterscheidet sich die RAM-Größe von GPU und CPU erheblich zugunsten der letzteren. Auf der Testplattform wurde 6 GB GDDR5 installiert, auf der CPU waren es sogar 48 GB. Die Speichergrenzen der GPU erlauben keine Berechnung großer Grafiken. Und hier können uns die CPU- und Unified Memory-Technologie helfen. [12]Hiermit können Sie auf die GPU im CPU-Speicher zugreifen. Da Informationen über das Diagramm nur für die Prozedur zum Ermitteln der Mindestkante erforderlich sind, können Sie bei großen Diagrammen Folgendes tun: Platzieren Sie zuerst alle Hilfsarrays im GPU-Speicher, und platzieren Sie dann einen Teil der Diagrammarrays (Array der Nachbarn A, Array X und Array der Gewichte W) im Speicher GPU, aber was hat nicht gepasst - in den Speicher der CPU. Ferner ist es während der Berechnung möglich, die Daten so zu unterteilen, dass der Teil, der nicht auf die GPU passte, auf der CPU verarbeitet wird und die GPU nur minimal auf den CPU-Speicher zugreift (da der Zugriff auf den CPU-Speicher vom Grafikbeschleuniger über den PCIe-Bus mit einer Geschwindigkeit von nicht mehr als 1 Sekunde erfolgt) 15 GB / s). Es ist im Voraus bekannt, in welchem ​​Verhältnis die Daten aufgeteilt wurden. Um zu bestimmen, auf welchen Speicher Sie in der GPU oder CPU zugreifen müssen, geben Sie einfach eine Konstante ein. Indem Sie anzeigen, an welcher Stelle die Arrays getrennt sind, und indem Sie den Algorithmus auf der GPU einmal überprüfen, können Sie bestimmen, wo Sie anrufen sollen. Die Position im Speicher dieser Arrays kann ungefähr so ​​dargestellt werden:

    Somit ist es möglich, Grafiken zu verarbeiten, die anfänglich nicht auf die GPU passen, selbst wenn die beschriebenen Komprimierungsalgorithmen verwendet werden, jedoch mit einer geringeren Geschwindigkeit, da der PCIe-Durchsatz sehr begrenzt ist.

    Testergebnisse


    Die Tests wurden mit der NVidia GTX Titan-GPU durchgeführt, die 14 SMX-Prozessoren mit 192 Cuda-Kernen (insgesamt 2688) und 6 Kernen (12.) Intel Xeon E5 v1660 mit einer Frequenz von 3,7 GHz aufweist. Die Grafiken, an denen Tests durchgeführt wurden, sind oben beschrieben. Ich werde nur einige Eigenschaften angeben:
    Skala (2 ^ N)Anzahl der Eckpunkte Anzahl der Rippen (2 * M)Diagrammgröße, GB
    RMAT SSCA2
    1665,5362 097 152 ~ 2 100 000 ~ 0,023
    212 097 152 67 108 864 ~ 67.200.000~ 0,760
    2416 777 216536 870 912 ~ 537 Millionen~ 6.3
    2533 554 4321,073,741,824 ~ 1.075.000.000~ 12,5
    2667 108 8642 147 483 648 ~ 2 150 000 000~ 25,2
    27134 217 7284,294,967,296 ~ 4.300.000.000~ 51,2

    Es ist zu erkennen, dass der Graph der Skala 16 recht klein ist (ca. 25 MB) und auch ohne Transformationen problemlos in den Cache eines modernen Intel Xeon-Prozessors passt. Und da die Diagrammgewichte 2/3 der Gesamtmenge einnehmen, müssen tatsächlich etwa 8 MB verarbeitet werden, was nur etwa dem Fünffachen des L2-GPU-Cache entspricht. Für große Grafiken ist jedoch ausreichend Speicher erforderlich, und selbst eine Grafik der Skala 24 kann ohne Komprimierung nicht mehr in den Speicher der getesteten GPU passen. Basierend auf der Graphendarstellung ist die 26. Skala die letzte, bei der die Anzahl der Kanten in ein vorzeichenloses int gesetzt wird, was eine Einschränkung des Algorithmus für die weitere Skalierung darstellt. Diese Einschränkung kann durch Erweitern des Datentyps leicht umgangen werden. Es scheint mir, dass dies bisher nicht so relevant ist, da die Verarbeitung einfacher Genauigkeit (unsigned int) um ein Vielfaches schneller ist. als double (unsigned long long) und die Speicherkapazität ist noch klein genug. Die Leistung wird in der Anzahl der Kanten gemessen, die pro Sekunde verarbeitet werden (überstrichene Kanten pro Sekunde - TEPS).

    Die Kompilierung wurde mit dem NVidia CUDA Toolkit 7.0 mit den Optionen -O3 -arch = sm_35 und Intel Composer 2015 mit den Optionen -O3 durchgeführt. Die maximale Leistung des implementierten Algorithmus ist in der folgenden Grafik dargestellt:

    Das Diagramm zeigt, dass bei Verwendung aller SSCA2-Optimierungen die Diagramme eine gute Leistung zeigen: Je größer das Diagramm, desto besser die Leistung. Dieses Wachstum wird beibehalten, bis alle Daten im Speicher der GPU abgelegt sind. Bei 25 und 26 Skalierungen wurde der Unified Memory-Mechanismus verwendet, der es ermöglichte, ein Ergebnis zu erzielen, wenn auch mit einer geringeren Geschwindigkeit (aber wie unten gezeigt wird, schneller als nur auf der CPU). Wenn die Berechnung auf einem Tesla k40 mit 12 GB Arbeitsspeicher und deaktiviertem ECC und Intel Xeon E5 V2 / V3-Prozessor durchgeführt würde, könnten auf einem SSCA2-Diagramm der Skala 25 etwa 3000 MTEPS erzielt und nicht nur ein Diagramm der 26. Skala verarbeitet werden, sondern auch 27. Ein solches Experiment wurde für den RMAT-Graphen aufgrund seiner komplexen Struktur und schlechten Anpassung des Algorithmus nicht durchgeführt.

    Vergleich der Leistung verschiedener Algorithmen


    Данная задача решалась в рамках конкурса конференции GraphHPC 2015. Я бы хотел привести сравнение с программой, написанной Александром Дарьиным, который по мнению авторов занял первое место в данном конкурсе.
    Так как в общей таблице есть результаты на тестовой платформе, предоставленной авторами, то не лишним было бы привести графики на CPU и GPU на описанной платформе (GTX Titan + Xeon E5 v2). Ниже представлены результаты для двух графов:



    Aus den Diagrammen geht hervor, dass der in diesem Artikel beschriebene Algorithmus für SSCA2-Diagramme optimiert ist, während der von Alexander Daryin implementierte Algorithmus für RMAT-Diagramme gut optimiert ist. In diesem Fall ist es unmöglich, eindeutig zu sagen, welche der Implementierungen die beste ist, da jede ihre eigenen Vor- und Nachteile hat. Auch das Kriterium, nach dem Algorithmen bewertet werden sollen, ist nicht klar. Wenn wir über die Verarbeitung großer Graphen sprechen, ist die Tatsache, dass der Algorithmus Graphen im Maßstab 24-26 verarbeiten kann, ein großes Plus und ein großer Vorteil. Wenn wir über die durchschnittliche Verarbeitungsgeschwindigkeit von Diagrammen jeder Größe sprechen, ist nicht klar, welcher Durchschnittswert zu berücksichtigen ist. Nur eines ist klar - ein Algorithmus verarbeitet SSCA2-Graphen gut, der zweite - RMAT. Wenn Sie diese beiden Implementierungen kombinieren, dann wird die durchschnittliche Leistung bei 23 Skalen etwa 3200 MTEPS betragen. Die Darstellung der Beschreibung einiger Optimierungen des Algorithmus von Alexander Darin kann eingesehen werdenhier .

    Von ausländischen Artikeln kann folgendes unterschieden werden.
    1) [13] Aus diesem Artikel wurden einige Ideen bei der Implementierung des beschriebenen Algorithmus verwendet. Es ist unmöglich, die von den Autoren erhaltenen Ergebnisse direkt zu vergleichen, da Tests mit dem alten NVidia Tesla S1070 durchgeführt wurden. Die von den Autoren auf der GPU erzielte Leistung reicht von 18 bis 36 MTEPS. Erschienen in 2009 und 2013.
    2) [14] Implementierung des Prim-Algorithmus auf der GPU.
    3) [15] Implementierung von k NN-Boruvka auf der GPU.
    Es gibt auch einige parallele Implementierungen auf der CPU. Aber ich konnte keine hohe Leistung in ausländischen Artikeln finden. Vielleicht kann einer der Leser sagen, ob ich etwas verpasst habe. Erwähnenswert ist auch, dass es in Russland praktisch keine Veröffentlichungen zu diesem Thema gibt (mit Ausnahme von Vadim Zaitsev ), was meines Erachtens sehr traurig ist.

    Über den Wettbewerb und statt eines Fazits


    Ich möchte meine Meinung über die Vergangenheit schreiben und den Wettbewerb um die beste Implementierung von MST erwähnen. Diese Notizen müssen nicht gelesen werden und geben meine persönliche Meinung wieder. Es ist möglich, dass jemand ganz anders denkt.

    Die Grundlage für die Lösung dieses Problems für alle Teilnehmer war derselbe Boruvka-Algorithmus. Es stellt sich heraus, dass die Aufgabe etwas vereinfacht wurde, da andere Algorithmen (Kruskala und Prima) einen hohen Rechenaufwand aufweisen und auf parallele Architekturen, einschließlich der GPU, langsam oder schlecht abgebildet sind. Aus dem Namen der Konferenz folgt logischerweise, dass ein Algorithmus geschrieben werden muss, der große Diagramme gut verarbeitet, z. B. Diagramme, die mindestens 1 GB Speicher belegen (solche Diagramme haben Skalen von 22 oder mehr). Leider haben die Autoren diese Tatsache aus irgendeinem Grund nicht berücksichtigt und die gesamte Konkurrenz musste einen Algorithmus schreiben, der im Cache gut funktioniert, da die Testplattform 2 CPUs mit einem Gesamt-Cache von 50 MB enthielt (Grafiken mit einem Gewicht von bis zu 17 Skalierungen <= 50 MB). Nur einer der Teilnehmer zeigte ein akzeptables Ergebnis - Vadim Zaitsev, einen ziemlich hohen Durchschnittswert für 2x CPU in Grafik 22 der Skala erhalten. Wie sich jedoch während der Konferenz herausstellte, ist dieser Teilnehmer seit geraumer Zeit mit der MST-Aufgabe befasst. Es ist wahrscheinlich, dass die Verarbeitungsgeschwindigkeit von großen Diagrammen anderer implementierter Algorithmen nicht groß ist und sich stark von den auf der Wettbewerbswebsite veröffentlichten Zahlen (zum Schlechten) unterscheidet. Zu beachten ist auch, dass die Graphenstrukturen sehr unterschiedlich sind und warum es plötzlich notwendig ist, den Durchschnittswert der Produktivität genau zu betrachten, da der arithmetische Durchschnitt ebenfalls nicht ganz klar ist. Die Größe des verarbeiteten Diagramms wurde ebenfalls nicht berücksichtigt. Ein weiteres unangenehmes "Feature" des bereitgestellten Systems (das 2x Intel Xeon E5-2690 und NVidia Tesla K20x enthielt) funktioniert nicht mit PCIe 3.0 (obwohl es von der GPU unterstützt wird und auf der Serverplatine vorhanden ist). Ergebend,

    Es sollte beachtet werden, dass das Lösen solcher Probleme auf der GPU nicht einfach ist, da die Graphverarbeitung auf der GPU aufgrund von Architekturmerkmalen schwierig zu parallelisieren ist. Und möglicherweise sollten diese Wettbewerbe zur Entwicklung von Spezialisten auf dem Gebiet der Schreibalgorithmen beitragen, die nicht-strukturelle Gitter für Grafikprozessoren verwenden. Nach den Ergebnissen dieses Jahres wie auch nach dem vorhergehenden Jahr ist der Einsatz der GPU bei solchen Aufgaben leider sehr begrenzt.

    Referenzen:


    [1] en.wikipedia.org/wiki/Sparse_matrix
    [2] www.dislab.org/GraphHPC-2014/rmat-siam04.pdf
    [3] www.dislab.org/GraphHPC-2015/SSCA2-TechReport.pdf
    [4 ] www.nvidia.ru/object/tesla-supercomputer-workstations-ru.html
    [5] www.nvidia.ru/object/geforce-gtx-titan-x-ru.html#pdpContent=2
    [6] www.amd .com / ru-ru / products / Grafik / Workstation / FirePro-3d / 9100
    [7] en.wikipedia.org/wiki/Bor%C5%AFvka ‚s_algorithm
    [8] www.cs.princeton.edu/~rs/ AlgsDS07 / 01UnionFind.pdf
    [9] habrahabr.ru/company/epam_systems/blog/247805
    [10]on-demand.gputechconf.com/gtc/2013/presentations/S3174-Kepler-Shuffle-Tips-Tricks.pdf
    [11] openmp.org/wp
    [12] devblogs.nvidia.com/parallelforall/unified-memory-in- cuda-6
    [13] stanford.edu/~vibhavv/papers/old/Vibhav09Fast.pdf
    [14] ieeexplore.ieee.org/xpl/login.jsp?tp=&arnumber=5678261&url=http%3A%2F%2Fieeexplore.ieee .org% 2Fxpls% 2Fabs_all.jsp% 3Farnumber% 3D5678261
    [15] link.springer.com/chapter/10.1007%2F978-3-642-31125-3_6

    Jetzt auch beliebt: