Lernen Sie Triton einen Kernel gleichzeitig: Vektor Zusatz

Eine kleine Optimierung reicht aus. Modelle mögen Gpt4 kosten mehr als 100 Millionen US -Greenback zum Coachingwas macht a 1% Effizienzgewinn Wert über eine Million Greenback. Eine leistungsstarke Möglichkeit, die Effizienz von Modellen für maschinelles Lernen zu optimieren, besteht darin, einige ihrer Komponenten zu schreiben direkt auf der GPU. Wenn Sie so etwas wie ich sind, reicht die einfache Erwähnung von Cuda -Kerneln aus, um Schüttelfrost in die Wirbelsäule zu schicken, da sie notorisch komplex sind, um zu schreiben und zu debuggen.

Glücklicherweise, Openai freigegeben Triton Im Jahr 2021, eine neue Sprache und Compiler, die einen Großteil der Komplexität von CUDA abstrahiert und weniger erfahrene Praktizierende ermöglicht, Performantenkörner zu schreiben. Ein bemerkenswertes Beispiel ist Nicht slothein LLM-Trainingsdienst, der verspricht 30x schnelleres Coaching mit 60% weniger Speicherverbrauchalles dank Ersetzen von Schichten in Pytorch durch Triton -Kernel.

In dieser Tutorial-Serie lernen wir die Grundlagen der GPU-Architektur und der Implementierung von Hochleistungs-Triton-Kerneln! Alle in dieser Serie vorgestellten Code sind bei verfügbar https://github.com/rpegoud/triton-kernels.

Grundlagen der GPU -Architektur

In diesem Abschnitt werden wir die Grundlagen von ((Nvidia) GPUs, um uns anzufangen und bis zum Ende dieses Artikels unseren ersten Triton -Kernel zu schreiben.

Ausgehend von der kleinsten Softwareinheit können wir die Hierarchie der Ausführungseinheiten wie folgt beschreiben:

• Themen: Der kleinste ArbeitseinheitSie führen den benutzerdefinierten Kernelcode aus.
• Verzerrt: Der kleinste PlanungseinheitSie bestehen immer aus 32 parallelen Threads mit jeweils ihren eigenen Anweisungsaditeln und Registrierungsstatus. Themen in einer Kette anfangen zusammen aber sind frei zu Zweig Und unabhängig ausführen.
• Fadenblöcke: Gruppe der Warps, wo alle Threads können kooperieren über gemeinsame Speicher und Synchronisierungsbarrieren. Es ist erforderlich, dass Threadblöcke ausgeführt werden können unabhängig und in jeder Reihenfolge parallel oder nacheinander. Diese Unabhängigkeit ermöglicht Threadblöcke in beliebiger Reihenfolge über eine beliebige Anzahl von Kernen geplantdamit GPU -Programme effizient mit der Anzahl der Kerne skalieren. Bei Bedarf können wir die Threads in einem Block an bestimmten Stellen im Kernel synchronisieren, um den Speicherzugriff zu synchronisieren.
• Streaming -Multiprozessor (SM): Eine Einheit, die verantwortlich ist viele Warps parallel ausführenEs besitzt den gemeinsamen Speicher und einen L1-Cache (enthält die neuesten globalen Reminiscence-Linien, auf die der SM zugegriffen hat). Ein SM hat eine engagierte Warp Scheduler Das Ziehen verzerrt aus den Fadenblöcken, die bereit sind, zu laufen.

Auf der Hardwareseite ist die kleinste Arbeitseinheit a Cuda -Kernder physische Arithmetische Logik UniT (Alu), der funktioniert arithmetische Operationen für einen Thread (oder Teile davon).

Um diesen Abschnitt mit einer Analogie zusammenzufassen, konnten wir sehen Cuda -Kerne als einzelne Arbeiterwährend a Kette ist a Truppe von 32 Arbeitern gleichzeitig die gleiche Anweisung. Sie können diese Aufgabe auf die gleiche Weise ausführen (verzweigen) und können sie möglicherweise zu einem anderen Zeitpunkt (Unabhängigkeit) ausführen. A Fadenblock besteht aus Mehrere Trupps teilen sich einen gemeinsamen Arbeitsbereich (dh das gemeinsame Gedächtnis) Arbeiter aus allen Trupps im Arbeitsbereich können aufeinander warten, um gleichzeitig zu Mittag zu essen. A Streaming -Multiprozessor ist a Fabrikboden mit vielen Trupps, die zusammenarbeiten und Instruments und Speicher teilen. Schließlich die GPU ist a ganze Pflanzemit vielen Böden.

Optimierungsgrundlagen

Bei der Optimierung von Deep Studying -Modellen jonglieren wir mit drei Hauptkomponenten:

1. Berechnen: Zeit, die von den GPU Computing Floating Level Operations (FLOPS) aufgewendet wird.
2. Erinnerung: Zeit für die Übertragung von Tensoren innerhalb einer GPU.
3. Überkopf: Alle anderen Operationen (Python -Dolmetscher, Pytorch Dispatch,…).

Wenn Sie diese Komponenten im Auge behalten, können Sie den richtigen Weg herausfinden, um einen Engpass zu lösen. Zum Beispiel hilft das Erhöhen des Berechnung (z. B. mit einer leistungsstärkeren GPU) nicht, wenn die meisten Zeit Speichertransfers ausgegeben werden. Im Idealfall sollte jedoch die meiste Zeit für Berechnung ausgegeben werden, genauer gesagt für Matrix -Multiplikationen, für die genaue Operation GPUs optimiert werden.

Dies impliziert, dass die Kosten für die Umschaltung von Daten entweder von der CPU in die GPU (”minimiert werden, um Daten umzusetzen (”Datenübertragungskosten”) Von einem Knoten zum anderen (”Netzwerkkosten”) Oder aus CUDA International Reminiscence (Drambillig aber langsam), um den gemeinsam genutzten Speicher (CUDA) (Sramteurer, aber schnellster Speicher des Geräts). Die später heißt es Bandbreite Kosten und wird vorerst unser Hauptaugenmerk sein. Zu den gängigen Strategien zur Reduzierung der Bandbreitenkosten gehören:

1. Wiederverwendung Daten in gemeinsamen Speicher geladen für mehrere Schritte. Ein Paradebeispiel hierfür ist die Multiplikation mit gekachelter Matrix, die wir in einem zukünftigen Beitrag behandeln werden.
2. Verschmelzen Mehrere Operationen in einem einzelnen Kernel (da jeder Kernel -Begin von DRAM zu SRAM bewegt wird), können wir beispielsweise eine Matrixmultiplikation mit einer Aktivierungsfunktion verschmelzen. Allgemein, Bedienerfusion Kann eine large Leistungserhöhung bieten, da es viele globale Speicherlesungen/Schreibvorgänge verhindert und zwei Betreiber eine Gelegenheit zur Fusion haben.

In diesem Beispiel führen wir eine Matrix -Multiplikation durch x@W und speichern das Ergebnis in einer Zwischenvariablen a. Wir wenden dann a an relu Zu a und speichern das Ergebnis in einer Variablen y. Dies erfordert, dass die GPU aus lesen kann x Und W Schreiben Sie das Ergebnis im globalen Gedächtnis in alesen aus a wieder und schließlich einschreiben y. Stattdessen würde uns Operator Fusion ermöglichen, die Anzahl der Lesevorgänge und Schreibvorgänge in International Reminiscence durch Durchführung der Matrix -Multiplikation und der Anwendung der Relu in einem einzelnen Kernel zu halbieren.

Triton

Wir werden jetzt unseren ersten Triton -Kernel schreiben, einen einfachen Vektor -Zusatz. Gehen wir zunächst durch, wie diese Operation auf einer GPU abgebaut und ausgeführt wird.

Erwägen Sie, die Einträge von zwei Vektoren zusammenzufassen wollen X Und Yjeweils mit 7 Elementen (n_elements=7).

Wir werden die GPU anweisen, dieses Downside in Stücken von 3 Elementen gleichzeitig anzugehen (BLOCK_SIZE=3). Um alle 7 Elemente der Eingangsvektoren abzudecken, startet die GPU 3 parallele „Programme“, unabhängige Instanz unseres Kernels, jeweils eine eindeutige Programm -ID, pid:

• Programm 0 wird Elemente zugewiesen 0, 1, 2.
• Programm 1 wird Elemente zugewiesen 3, 4, 5.
• Programm 2 wird Component zugewiesen 6.

Dann schreiben diese Programme die Ergebnisse in einem Vektor zurück Z in globalem Gedächtnis gespeichert.

Ein wichtiges Element ist, dass ein Kernel keinen ganzen Vektor erhält Xstattdessen erhält es a Zeiger auf die Speicheradresse des ersten ComponentsAnwesend X(0). Um auf die tatsächlichen Werte von zuzugreifen XWir müssen sie manuell aus dem globalen Speicher laden.

Wir können auf die Daten für jeden Block zugreifen, indem wir die Programm -ID verwenden: block_start = pid * BLOCK_SIZE. Von dort aus können wir die verbleibenden Elementadressen für diesen Block durch Pc erhalten offsets = block_start + vary(0, BLOCK_SIZE) und laden Sie sie in den Speicher.

Denken Sie jedoch daran, dass Programm 2 nur Component 6 zugewiesen ist, aber seine Offsets sind es (6, 7, 8). Um einen Indexierungsfehler zu vermeiden, lässt uns Triton a definieren Maske Hier zu identifizieren, um gültige Zielelemente zu identifizieren masks = offsets < n_elements.

Wir können jetzt sicher laden X Und Y und fügen Sie sie zusammen, bevor Sie das Ergebnis zu einer Ausgangsvariablen zurückschreiben Z in globalem Gedächtnis in ähnlicher Weise.

Schauen wir uns den Code genauer an, hier ist der Triton -Kernel:

import triton
import triton.language as tl

@triton.jit
def add_kernel(
	x_ptr, # pointer to the primary reminiscence entry of x
	y_ptr, # pointer to the primary reminiscence entry of y
	output_ptr, # pointer to the primary reminiscence entry of the output
	n_elements, # dimension of x and y
	BLOCK_SIZE: tl.constexpr, # dimension of a single block
):
	# --- Compute offsets and masks ---
	pid = tl.program_id(axis=0) # block index
	block_start = pid * BLOCK_SIZE # begin index for present block
	offsets = block_start + tl.arange(0, BLOCK_SIZE) # index vary
	masks = offsets < n_elements # masks out-of-bound parts
	
	# --- Load variables from world reminiscence ---
	x = tl.load(x_ptr + offsets, masks=masks)
	y = tl.load(y_ptr + offsets, masks=masks)

	# --- Operation ---
	output = x + y	
	
	# --- Save outcomes to world reminiscence ---
	tl.retailer(pointer=output_ptr + offsets, worth=output, masks=masks)

Lassen Sie uns einige der tritonspezifischen Syntax aufschlüsseln:

• Erstens ist ein Triton -Kernel immer von dekoriert von <a href="http://twitter.com/triton" goal="_blank" rel="noreferrer noopener">@triton</a>.jit.
• Zweitens müssen einige Argumente als statisch deklariert werden, was bedeutet, dass sie zur Rechenzeit bekannt sind. Dies ist erforderlich für BLOCK_SIZE und wird erreicht, indem das hinzufügen tl.constexpr Geben Sie Annotation ein. Beachten Sie auch, dass wir keine anderen Variablen kommentieren, da sie keine ordnungsgemäßen Python -Variablen sind.
• Wir verwenden tl.program_id Zugriff auf die ID des aktuellen Blocks, tl.arange verhält sich ähnlich wie Numpy’s np.arange.
• Das Laden und Speichern von Variablen wird durch Anrufe erreicht tl.load Und tl.retailer mit Anordnungen von Zeigern. Beachten Sie, dass es keine gibt return Aussage, diese Rolle wird delegiert tl.retailer.

Um unseren Kernel zu benutzen, müssen wir jetzt a schreiben Wrapper auf Pytorch-Ebene Das liefert Speicherzeiger und definiert a Kernelgitter. Im Allgemeinen ist das Kernelnetz ein 1D-, 2D- oder 3D -Tupel, das das enthält Anzahl der am Kernel zugewiesenen Fadenblöcken entlang jeder Achse. In unserem vorherigen Beispiel haben wir ein 1D -Gitter mit 3 Fadenblöcken verwendet: grid = (3, ).

Um unterschiedliche Arraygrößen zu behandeln, sind wir standardmäßig grid = (ceil(n_elements / BLOCK_SIZE), ).

def add(X: torch.Tensor, Y: torch.Tensor) -> torch.Tensor:
	"""PyTorch wrapper for `add_kernel`."""
	output = torch.zeros_like(x) # allocate reminiscence for the output
	n_elements = output.numel()  # dimension of X and Y
	
	# cdiv = ceil div, computes the variety of blocks to make use of
	grid = lambda meta: (triton.cdiv(n_elements, meta("BLOCK_SIZE")),)
	# calling the kernel will robotically retailer `BLOCK_SIZE` in `meta`
	# and replace `output`
	add_kernel(grid)(X, Y, output, n_elements, BLOCK_SIZE=1024)
	
	return output

Hier sind zwei letzte Notizen über den Wrapper:

Sie haben das vielleicht bemerkt grid wird als Lambda -Funktion definiert. Dadurch kann Triton die Anzahl der zu starteten Threadblöcke berechnen zur Startzeit. Daher berechnen wir die Gittergröße basierend auf der in gespeicherten Blockgröße metaein Wörterbuch von Kompilierungs-Zeitkonstanten, die dem Kernel ausgesetzt sind.

Wenn Sie den Kernel nennen, den Wert von output wird an Ort und Stelle geändert, sodass wir uns nicht neu zuweisen müssen output = add_kernel(…).
Wir können dieses Tutorial abschließen, indem wir überprüfen, ob unser Kernel ordnungsgemäß funktioniert:

x, y = torch.randn((2, 2048), machine="cuda")

print(add(x, y))
>> tensor(( 1.8022, 0.6780, 2.8261, ..., 1.5445, 0.2563, -0.1846), machine='cuda:0')

abs_difference = torch.abs((x + y) - add(x, y))
print(f"Max absolute distinction: {torch.max(abs_difference)}")
>> Max absolute distinction: 0.0

Das ist es für diese Einführung. In folgenden Beiträgen werden wir lernen, interessantere Kernel wie geflieste Matrix -Multiplikation zu implementieren und zu sehen autograd.

Bis zum nächsten Mal! 👋

Referenzen und nützliche Ressourcen