AMD ROCm: 4-Wave-Interleave-Optimierung der FP8-GEMM-Kernel fuer Instinct MI355X (CDNA 4) verdoppelt Register-Budget pro Wave

Das AMD ROCm Team veroeffentlichte eine detaillierte technische Beschreibung einer neuen Optimierungsstrategie fuer FP8-GEMM-Kernel auf AMD Instinct MI355X GPUs mit CDNA-4-Architektur. Die Optimierung geht ueber den Standard-8-Wave-Ping-Pong-Ansatz hinaus und fuehrt ein 4-Wave-Interleave-Design ein, das das Register-Budget pro Wave deutlich erhoeht und Speicherkonflikte auf Hardware-Ebene eliminiert.

Warum dominiert der 4-Wave-Ansatz fuer FP8 GEMM?

Der wesentliche Unterschied liegt in der VGPR-Zuteilung (Vector General Purpose Registers). Im 8-Wave-Design teilt jede Wave Registerplatz, was die Output-Tile-Groesse begrenzt. Der 4-Wave-Ansatz weist einer Wave pro SIMD-Einheit zu und gibt ihr Zugriff auf das volle Budget von 512 VGPR-Registern.

Dies wirkt sich direkt auf die Tile-Groesse aus: 4-Wave ermoeglicht 128x128-Output-Tiles gegenueber 64x128 beim 8-Wave-Design. Der globale Tile betraegt 256x256x128, und jede der vier Waves verarbeitet vier 64x64-Bloecke mit 16x16-MFMA-Operationen.

Wie eliminiert 4-Wave-Interleave Speicherlatenz auf dem MI355X?

Die zentrale Herausforderung bei GPU-Matrixmultiplikations-Kerneln ist die Latenzverbergung. Der 4-Wave-Ansatz loest dies durch praezises Verflechten von Instruktionen auf Sub-Tile-Granularitaet. Die Implementierung nutzt __builtin_amdgcn_sched_barrier(0). Jeder interleaved_block-Aufruf ueberlappt 16 MFMA-Instruktionen mit 8 LDS-zu-Register-Ladevorgaengen ohne Ping-Pong-Synchronisierungskomplexitaet.

Wie eliminiert XOR-basiertes LDS-Swizzle Speicherbankkonflikte?

Der MI355X-LDS-Speicher hat 64 Baenke a 4 Byte. 16x16x128-MFMA-Instruktionen greifen naturgemaeass spaltenweise zu, was 8-fache Bankkonflikte verursacht. AMD loest dies mit:

row_bits = (offset % (16 * 128) >> 7) / 2
mask = row_bits << 4
swizzled_offset = offset ^ mask

Die XOR-Operation remappt Adressen fuer konfliktfreien Ablauf aller vier ds_read_b128-Phasen — voller LDS-Durchsatz statt effektiver 12,5% bei 8-fachen Konflikten.

Hardware-Plattform und Chiplet-aware-Optimierung

Alle Optimierungen zielen auf den AMD Instinct MI355X mit CDNA-4-Architektur und 8-Chiplet-(XCD)-Topologie. Das ROCm-Team ergaenzt ein Chiplet-aware Grid Swizzle mit Morton-Curve-aehnlichem Mapping fuer maximale Cache-Nutzung. Benchmarking erfolgte mit 1000 Warm-up- und 1000 Benchmark-Iterationen; das Team meldet konsistente Leistung ueber ROCm-Versionen ohne manuelles Pragma-Tuning.