In dit artikel beargumenteren de auteurs dat traditionele abstractielagen in softwareontwikkeling, zoals CUDA DSL's, op weg zijn naar hun pensioen. Waar voorheen complexe C++-abstracties essentieel waren om de cognitieve belasting bij het schrijven van 'megakernels' te beheersen, kunnen AI-agents nu direct target-geoptimaliseerde code genereren vanuit vage instructies.
De kern van hun betoog is dat de rol van abstractie als 'cognitieve ontlaster' verdwijnt. Er ontstaat een verschuiving van een focus op precisie en fragiele codebases naar flexibele prompts die intenties overbrengen. Hoewel de auteurs erkennen dat er uitdagingen blijven rondom verificatie (het belang van een 'orakel' of referentie-implementatie) en kennisoverdracht, concluderen zij dat de implementatie zelf wegwerpbaar wordt.
Uiteindelijk verschuift het vertrouwen van de specifieke code ('de diff') naar de specificaties, invarianten en domeinkennis. De codebase is niet langer de bron van waarheid; de ideeën erachter blijven dat wel.
De abstracties met pensioen sturen
TL;DR: We denken dat CUDA DSL's op weg zijn naar hun pensioen.
Vorig jaar zetten we ons af op het schrijven van megakernels. Maar het implementeren van megakernels kan ronduit ellendig zijn. Het vereist complexe datastructuren, uitgebreide synchronisatie tussen threads/SM/GPU, diep geneste control flow, enzovoort. We konden dit niet allemaal in ons hoofd houden, dus deden we wat de informatica al zeventig jaar doet: we bouwden een abstractielaag. Zelfs daarmee moesten we enkele maanden vechten tegen race-conditions en deadlocks voordat de Llama's razendsnel draaiden.
Dit jaar hebben we een MoE-megakernel gebouwd, maar we hebben de abstractie verwijderd. Met behulp van agents konden we direct met de complexiteit werken en vanaf nul target-geoptimaliseerde code bouwen. Tussenliggende lagen van C++-abstracties waren niet nodig.
Wat zou dit kunnen betekenen?
De ontlaster ontlasten
We hebben het afgelopen kwartaal een interessant patroon opgemerkt. Taken die we al zonder abstracties konden uitvoeren (bijvoorbeeld het schrijven van een geoptimaliseerde GEMM-kernel), zijn nu bijna volledig geautomatiseerd mits de juiste prompt wordt gegeven. Voorlopig moeten we de agent nog steeds vertellen welke PTX-instructies hij moet gebruiken, wat het ontwerp voor warp-specialisatie is, enzovoort, maar hij bereikt de state-of-the-art zeer snel.
Taken waarvoor we abstracties nodig hadden (bijvoorbeeld het schrijven van een megakernel), zijn helaas nog niet geautomatiseerd; we kunnen vandaag de dag zeker geen megakernel in één prompt ('one-shot') genereren. Maar agents stellen ons in staat om de abstractie in de prompt te plaatsen — in een onvolledige, rommelige vorm — in plaats van in zorgvuldig ontworpen C++ templates. Complexiteit die vorig jaar onbeheersbaar was, is plotseling beheersbaar geworden, omdat we nu een "compiler" hebben die vage instructies kan omzetten in code. De ideeën die ten grondslag lagen aan de megakernel bleven net zo belangrijk, maar ze waren veel gemakkelijker uit te drukken.
De rol van abstractie als cognitieve ontlaster begint te verdwijnen. Agents nemen dat werk over.
Wat volgt er nu?
Door inductief redeneren zijn CUDA DSL's de volgende op de lijst voor het pensioen, inclusief onze geliefde ThunderKittens¹. Waarschijnlijk volgend jaar, of misschien zelfs eerder. We weten niet hoe ver dit zal gaan. Ergens onder ons ligt een bodem waar een abstractie puur het contract is, in plaats van een ontlaster, en daar stopt het. Maar we kunnen die bodem vanaf hier nog niet zien.
Dit leidt tot de provocatie: als een prompt elke laag van de stack kan genereren, wat is een codebase dan nog waard?
De kracht van een codebase is precisie. Het is het enige artefact dat niet ambigu is; het ding dat een specifieke machine twee keer op precies dezelfde manier zal uitvoeren. Maar precisie heeft een prijs, en die prijs is fragiliteit. Een codebase is gebonden aan een taal, een framework, een hardware-target en een set conventies die alleen het team dat het geschreven heeft echt begrijpt. Het lijkt draagbaar (portable), maar is dat nauwelijks.
Een prompt is het tegenovergestelde. Het is vaag, maar het is verplaatsbaar. Geef dezelfde intentie aan een andere werker, mens of machine, en als die werker intelligent genoeg is om de gaten correct in te vullen, krijg je een correct resultaat zonder dat de gaten vooraf gestandaardiseerd hoeven te zijn.
Dat is het interessante deel. Met agents is er nu een concurrentiestrijd tussen een intelligente uitvoerder die ondergespecificeerde instructies leest en het juiste doet, versus het specificeren van elk detail zodat een domme uitvoerder geen oordeel nodig heeft. DSL's en frameworks zijn die specificatie, uitgedrukt als een codebase. Als de uitvoerder niet langer dom is, verliest de DSL zijn bestaansrecht.
Een codebase is een imperfect medium. We houden ervan, maar het is enorm, barok en vol bugs. Niet omdat de ingenieurs die het schreven onvoorzichtig waren, maar omdat de volgende set beperkingen bijna onmogelijk is voor een reeks tokens om te vervullen:
- precies genoeg voor een machine;
- leesbaar genoeg voor een mens;
- flexibel genoeg voor de toekomst.
En we hebben eerder ondoorzichtige transformaties geaccepteerd. Weinigen schrijven de assembly die daadwerkelijk wordt uitgevoerd. De compiler herordent, inlined, vectoriseert en wijzigt wat er wordt uitgevoerd op manieren waar niemand in het team naar kan wijzen. We gingen akkoord met die deal omdat de compiler vaker gelijk had dan wij. Intelligentie zit al tussen de intentie en de machine. De vraag is of het toestaan dat deze intelligentie hoger in de stack zit, een verschil in graad of een verschil in soort is; wij denken dat het geen verschil in soort is.
Toelaatbaarheid
Toegeven, we zijn wat kort door de bocht geweest over de functie van een abstractie. Er is meer te overwegen voordat we er een loslaten.
Ten eerste is een abstractie niet alleen een manier om cognitieve belasting te beheersen; het is een gedeeld oppervlak waar applicatie, hergebruik en review aan vastklikken. Zonder dit krijg je een "Cambrische explosie" van verificatie-uitdagingen. Tien teams die ThunderKittens gebruiken, testen dezelfde tile-semantiek, en dat cumuleert in de loop van de tijd. Tien teams die op maat gemaakte megakernels genereren, hebben tien verschillende sets problemen die niet op elkaar voortbouwen.
Ten tweede is er geen gedeeld orakel. Voor een megakernel zijn de tests het contract én het werk. Wie schrijft ze, en wie controleert die tests? Toen we het framework verwijderden, behielden we het orakel: referentie-implementaties, numerieke toleranties, onze intuïtie over hoe het geprofileerde Gantt-diagram eruit zou moeten zien, enzovoort. Dat is wat de gepensioneerde abstractie heeft achtergelaten. We kunnen een laag pas pensioneren als we een orakel hebben dat langer meegaat dan die laag. In domeinen waar niemand weet wat het orakel is, is dit niet van toepassing en zouden we de steigers behouden.
Ten slotte zijn wij een enkel, bevooroordeeld voorbeeld. We hebben abstracties verwijderd in een domein dat we diepgaand kenden. Abstracties bestaan ook om kennis over te dragen aan mensen die die kennis nog niet hebben. We kunnen niet zeggen of een agent die de gaten voor ons invult, dit ook doet voor iemand die voor het eerst hoort over warp-specialisatie. Misschien maken agents onboarding-lagen overbodig. Of misschien hebben mensen ze nog steeds nodig, en hebben we simpelweg beschreven hoe dat is vanuit een comfortabele positie voor mensen die al getraind zijn.
Het pensioen
Ongeacht deze punten lijkt de transformatie onvermijdelijk, en we moeten onze geliefde frameworks wellicht loslaten.
Wat houden we dan over? De intentie, de invarianten, de tests en de zwaarbevochten domeinkennis die momenteel in de abstracties van ThunderKittens leeft en ze correct houdt op hardware. De bibliotheek mag verdwijnen. Maar die kennis verdampt niet, alleen omdat de onderliggende CUDA of HIP opnieuw wordt gegenereerd in plaats van handmatig wordt afgesteld.
En wat verandert er? Vertrouwen verschuift een niveau omhoog. We controleren de specificatie en het orakel, niet de 'diff'. De implementatie wordt wegwerpbaar: een cache van één specifieke compilatie, niet de bron van waarheid. Codebases zullen niet verdwijnen, maar ze stoppen met zijn wat we bedoelen als we spreken over "het project".
Abstracties gaan met pensioen. Ideeën blijven.
***
¹ We zullen onze ThunderKittens nog steeds updaten, onderhouden en koesteren. Vera Rubin-kernels volgen binnenkort!
De abstracties met pensioen sturen
TL;DR: We denken dat CUDA DSL's op weg zijn naar hun pensioen.
Vorig jaar zetten we ons af op het schrijven van megakernels. Maar het implementeren van megakernels kan ronduit ellendig zijn. Het vereist complexe datastructuren, uitgebreide synchronisatie tussen threads/SM/GPU, diep geneste control flow, enzovoort. We konden dit niet allemaal in ons hoofd houden, dus deden we wat de informatica al zeventig jaar doet: we bouwden een abstractielaag. Zelfs daarmee moesten we enkele maanden vechten tegen race-conditions en deadlocks voordat de Llama's razendsnel draaiden.
Dit jaar hebben we een MoE-megakernel gebouwd, maar we hebben de abstractie verwijderd. Met behulp van agents konden we direct met de complexiteit werken en vanaf nul target-geoptimaliseerde code bouwen. Tussenliggende lagen van C++-abstracties waren niet nodig.
Wat zou dit kunnen betekenen?
De ontlaster ontlasten
We hebben het afgelopen kwartaal een interessant patroon opgemerkt. Taken die we al zonder abstracties konden uitvoeren (bijvoorbeeld het schrijven van een geoptimaliseerde GEMM-kernel), zijn nu bijna volledig geautomatiseerd mits de juiste prompt wordt gegeven. Voorlopig moeten we de agent nog steeds vertellen welke PTX-instructies hij moet gebruiken, wat het ontwerp voor warp-specialisatie is, enzovoort, maar hij bereikt de state-of-the-art zeer snel.
Taken waarvoor we abstracties nodig hadden (bijvoorbeeld het schrijven van een megakernel), zijn helaas nog niet geautomatiseerd; we kunnen vandaag de dag zeker geen megakernel in één prompt ('one-shot') genereren. Maar agents stellen ons in staat om de abstractie in de prompt te plaatsen — in een onvolledige, rommelige vorm — in plaats van in zorgvuldig ontworpen C++ templates. Complexiteit die vorig jaar onbeheersbaar was, is plotseling beheersbaar geworden, omdat we nu een "compiler" hebben die vage instructies kan omzetten in code. De ideeën die ten grondslag lagen aan de megakernel bleven net zo belangrijk, maar ze waren veel gemakkelijker uit te drukken.
De rol van abstractie als cognitieve ontlaster begint te verdwijnen. Agents nemen dat werk over.
Wat volgt er nu?
Door inductief redeneren zijn CUDA DSL's de volgende op de lijst voor het pensioen, inclusief onze geliefde ThunderKittens¹. Waarschijnlijk volgend jaar, of misschien zelfs eerder. We weten niet hoe ver dit zal gaan. Ergens onder ons ligt een bodem waar een abstractie puur het contract is, in plaats van een ontlaster, en daar stopt het. Maar we kunnen die bodem vanaf hier nog niet zien.
Dit leidt tot de provocatie: als een prompt elke laag van de stack kan genereren, wat is een codebase dan nog waard?
De kracht van een codebase is precisie. Het is het enige artefact dat niet ambigu is; het ding dat een specifieke machine twee keer op precies dezelfde manier zal uitvoeren. Maar precisie heeft een prijs, en die prijs is fragiliteit. Een codebase is gebonden aan een taal, een framework, een hardware-target en een set conventies die alleen het team dat het geschreven heeft echt begrijpt. Het lijkt draagbaar (portable), maar is dat nauwelijks.
Een prompt is het tegenovergestelde. Het is vaag, maar het is verplaatsbaar. Geef dezelfde intentie aan een andere werker, mens of machine, en als die werker intelligent genoeg is om de gaten correct in te vullen, krijg je een correct resultaat zonder dat de gaten vooraf gestandaardiseerd hoeven te zijn.
Dat is het interessante deel. Met agents is er nu een concurrentiestrijd tussen een intelligente uitvoerder die ondergespecificeerde instructies leest en het juiste doet, versus het specificeren van elk detail zodat een domme uitvoerder geen oordeel nodig heeft. DSL's en frameworks zijn die specificatie, uitgedrukt als een codebase. Als de uitvoerder niet langer dom is, verliest de DSL zijn bestaansrecht.
Een codebase is een imperfect medium. We houden ervan, maar het is enorm, barok en vol bugs. Niet omdat de ingenieurs die het schreven onvoorzichtig waren, maar omdat de volgende set beperkingen bijna onmogelijk is voor een reeks tokens om te vervullen:
- precies genoeg voor een machine;
- leesbaar genoeg voor een mens;
- flexibel genoeg voor de toekomst.
En we hebben eerder ondoorzichtige transformaties geaccepteerd. Weinigen schrijven de assembly die daadwerkelijk wordt uitgevoerd. De compiler herordent, inlined, vectoriseert en wijzigt wat er wordt uitgevoerd op manieren waar niemand in het team naar kan wijzen. We gingen akkoord met die deal omdat de compiler vaker gelijk had dan wij. Intelligentie zit al tussen de intentie en de machine. De vraag is of het toestaan dat deze intelligentie hoger in de stack zit, een verschil in graad of een verschil in soort is; wij denken dat het geen verschil in soort is.
Toelaatbaarheid
Toegeven, we zijn wat kort door de bocht geweest over de functie van een abstractie. Er is meer te overwegen voordat we er een loslaten.
Ten eerste is een abstractie niet alleen een manier om cognitieve belasting te beheersen; het is een gedeeld oppervlak waar applicatie, hergebruik en review aan vastklikken. Zonder dit krijg je een "Cambrische explosie" van verificatie-uitdagingen. Tien teams die ThunderKittens gebruiken, testen dezelfde tile-semantiek, en dat cumuleert in de loop van de tijd. Tien teams die op maat gemaakte megakernels genereren, hebben tien verschillende sets problemen die niet op elkaar voortbouwen.
Ten tweede is er geen gedeeld orakel. Voor een megakernel zijn de tests het contract én het werk. Wie schrijft ze, en wie controleert die tests? Toen we het framework verwijderden, behielden we het orakel: referentie-implementaties, numerieke toleranties, onze intuïtie over hoe het geprofileerde Gantt-diagram eruit zou moeten zien, enzovoort. Dat is wat de gepensioneerde abstractie heeft achtergelaten. We kunnen een laag pas pensioneren als we een orakel hebben dat langer meegaat dan die laag. In domeinen waar niemand weet wat het orakel is, is dit niet van toepassing en zouden we de steigers behouden.
Ten slotte zijn wij een enkel, bevooroordeeld voorbeeld. We hebben abstracties verwijderd in een domein dat we diepgaand kenden. Abstracties bestaan ook om kennis over te dragen aan mensen die die kennis nog niet hebben. We kunnen niet zeggen of een agent die de gaten voor ons invult, dit ook doet voor iemand die voor het eerst hoort over warp-specialisatie. Misschien maken agents onboarding-lagen overbodig. Of misschien hebben mensen ze nog steeds nodig, en hebben we simpelweg beschreven hoe dat is vanuit een comfortabele positie voor mensen die al getraind zijn.
Het pensioen
Ongeacht deze punten lijkt de transformatie onvermijdelijk, en we moeten onze geliefde frameworks wellicht loslaten.
Wat houden we dan over? De intentie, de invarianten, de tests en de zwaarbevochten domeinkennis die momenteel in de abstracties van ThunderKittens leeft en ze correct houdt op hardware. De bibliotheek mag verdwijnen. Maar die kennis verdampt niet, alleen omdat de onderliggende CUDA of HIP opnieuw wordt gegenereerd in plaats van handmatig wordt afgesteld.
En wat verandert er? Vertrouwen verschuift een niveau omhoog. We controleren de specificatie en het orakel, niet de 'diff'. De implementatie wordt wegwerpbaar: een cache van één specifieke compilatie, niet de bron van waarheid. Codebases zullen niet verdwijnen, maar ze stoppen met zijn wat we bedoelen als we spreken over "het project".
Abstracties gaan met pensioen. Ideeën blijven.
***
¹ We zullen onze ThunderKittens nog steeds updaten, onderhouden en koesteren. Vera Rubin-kernels volgen binnenkort!
Groetjes,