Unser vorheriger Beitrag verfolgte einen Vektoradditions-Kernel von nvcc bis zu den Warps. Diesmal verfolgen wir eine kritische SASS-Instruktion (eine globale Ladung) durch die Hardware einer RTX4090. Der untersuchte CUDA-Kernel hat zwei Zeilen: __global__ void vadd(const float* a, const float* b, float* c, int n) {int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i];}. Das kompilierte SASS zeigt: /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i]. Diese Instruktionen laden Elemente des Vektors b aus dem globalen Speicher in die Register.
Eine LDG.E fordert vier Bytes in jedem der 32 Lanes an. Ihre Bedienung erfordert vier 32-Byte-Sektoren, eine Cache-Zeile, eine Adressübersetzung, ein Kreuzen des Crossbars, einen der sechsunddreißig L2-Slices und, wenn überall ein Fehlschlag vorliegt, eine Aktivierung und vier Spaltenlesevorgänge auf einem DRAM-Chip. Unser Warp lebt auf einer der vier Unterteilungen des SM neben elf anderen residenten Warps. Die Instruktion wird an die Load/Store-Einheit (LSU) ausgegeben, die den Opcode, die aktive Lane-Maske, die berechneten Adressen und die Registernummer an den Coalescer sendet.
Der Coalescer stellt fest, dass für die 128 Bytes, die der Warp angefordert hat, 4 zusammenhängende Sektoranforderungen benötigt werden. Beim Eintritt in den L1-Cache wird die Anforderung an den virtuell adressierten, 4-fach satzassoziativen Cache gesendet. Die Suche vergleicht Tags in Sätzen von vier Slots. Da dies eine erstmalige Ladung ist, verfehlt sie und steigt weiter hinab. Ein L1-Treffer kehrt in etwa 15,4 ns zurück — 40 Zyklen. Virtueller Speicher erfordert eine Übersetzung, bevor er den SM verlässt. Die TLB des SM enthält sechzehn aktuelle Übersetzungen. Fehlschläge kosten etwa 4,4 ns — elf Zyklen.
Nach der Übersetzung geht eine Anforderung pro 128-Byte-Zeile über den Crossbar zu einem der 36 L2-Slices mit 2 MiB, ausgewählt durch eine Funktion seiner physikalischen Adresse. Alle Slices können parallel bedienen, was eine aggregierte Bandbreite ergibt, die 36-mal so hoch ist wie die eines einzelnen Slices.
Quelle: Hacker News · Zusammengefasst von HeadlinesBriefing