Über die Kachelphasen iterieren

Eine vollständige Zeile mal eine vollständige Spalte ist zu groß für ein Tile. Daher teilen Sie die lange Summe in Abschnitte der Breite TILE auf – einen Abschnitt pro Phase.

Sind die Matrizen N Elemente breit und die Tiles TILE Elemente breit, benötigen Sie N / TILE Phasen, um die gesamte innere Dimension abzudecken.

Fassen Sie die Schritte Laden, Synchronisieren und Berechnen in einer Schleife über phase zusammen. Jeder Durchlauf verschiebt das Tile-Fenster weiter entlang der Zeile von A und der Spalte von B.

Die lokale Variable sum befindet sich außerhalb der Schleife und wächst daher weiter. Jede Phase addiert ihren Anteil am Skalarprodukt zur laufenden Summe.

Jede Phase verschiebt die aus A gelesene Spalte und die aus B gelesene Zeile um phase * TILE. So bewegt sich das Fenster weiter.

Rufen Sie wie zuvor nach dem Laden __syncthreads() auf, damit das Tile vollständig ist, bevor jemand damit rechnet.

Die innere Schleife addiert TILE Produkte zu sum und verwendet dabei ausschließlich das frisch geladene Tile. Sie liefert einen Abschnitt des endgültigen Skalarprodukts.

Beenden Sie jede Phase mit einem weiteren __syncthreads(), damit kein Thread das Tile überschreibt, während ein langsamerer Thread noch daraus liest. 🚧

Eine Barriere schützt Lesezugriffe nach dem Laden, die andere schützt Ladevorgänge nach dem Lesen. Gemeinsam halten sie jeden Thread über die Phasen hinweg im Gleichschritt.

Nach Abschluss aller Phasen ist die laufende sum das vollständige Skalarprodukt. Speichern Sie sie genau einmal in C und sichern Sie den Zugriff durch eine Bereichsprüfung ab.

Wenn N kein ganzzahliges Vielfaches von TILE ist, laden Sie für Elemente außerhalb des gültigen Bereichs zero, damit die zusätzlichen Produkte nichts zur Summe beitragen.

Sie haben über phases iteriert, Tiles verschoben, zweimal synchronisiert und Teilsummen zu einer Gesamtsumme akkumuliert. Nun messen wir, wie viel schneller das ist. 📈