Triton 3.8.0 překládá jádra pro Rubin, který je zatím jen v náhledu dokumentace PTX
Vydání z 28. srpna přidalo do překladače Triton cíl „cuda:107“ pro architekturu, kterou NVIDIA popisuje zatím jen ve vývojářském náhledu dokumentace PTX. Z Tritonu i z toho náhledu se dá spočítat, kolik má Rubin tenzorové paměti: 576 sloupců proti 512 na Blackwellu.
Triton je překladač, ve kterém se jádra pro grafické karty píšou v Pythonu místo v CUDA C. Sám o sobě žádný model nespustí; je to vrstva pod tím, co jádra potřebuje – balíček torch si ho na Linuxu instaluje jako závislost. Co Triton neumí přeložit, to na daném čipu nepoběží.
Verze 3.8.0 vyšla na GitHubu 28. srpna 2026 v 18:25 UTC, kola na PyPI byla nahraná o dvě hodiny dřív. Předchozí stabilní vydání 3.7.1 je z 18. června. V poznámkách stojí věta, kvůli které se to vydání vyplatí otevřít: „Added initial NVIDIA Rubin (SM107) support.“
Odkazuje na tři návrhy změn a hlavní z nich, číslo 10936, byl přijatý už 17. července: 27 souborů, 902 přidaných řádků. Do verze 3.7.1 se nedostal. Soubor s jeho testy, test/Conversion/tritongpu_to_llvm_rubin.mlir, na značce v3.7.1 neexistuje, na v3.8.0 ano. V mezijazyce se nový cíl píše jako ttg.target = "cuda:107".
Tenzorová paměť má o osminu víc sloupců
Pátá generace tenzorových jader má vlastní paměť přímo na čipu. Dokumentace PTX jí říká Tensor Memory a popisuje ji jako mřížku o 128 řádcích, kde každá buňka drží 32 bitů; na Blackwellu má 512 sloupců. Tabulka 58 v náhledu dokumentace uvádí u výlučné alokace strop 512 sloupců pro sm_100, sm_103 i sm_110, ale 576 pro sm_107. Triton má totéž číslo v hlavičce TargetFeatures.h: metoda getMaxTMEMColumns() vrátí 576, když se výpočetní schopnost rovná 107, a jinak 512.
Přepočet na bajty je náš a je jednoduchý: 576 sloupců krát 128 řádků krát čtyři bajty dá 294 912 B, kdežto 512 krát 128 krát čtyři jen 262 144 B. Rozdíl 32 kB sedí na to, co autor návrhu napsal do jeho popisu – „TMEM capacity per SM to 288KB“. K větší paměti přibyl v PTX i příznak .exclusive u instrukcí tcgen05.alloc a tcgen05.dealloc: výlučná alokace obvykle zabere celou tenzorovou paměť a žádná jiná vedle ní žít nesmí.
Sdílená paměť nad 228 kB se platí cache
Druhá stopa je v ovladači, kterým Triton jádra spouští. Do souboru driver.c přibyla kontrola, že přihlašovaná sdílená paměť nepřesáhne 228 kB – tolik jí má na multiprocesor Blackwell –, a za ní větev pro případ, že jádro chce víc. Tehdy spuštění dostane atribut CU_LAUNCH_ATTRIBUTE_SHARED_MEMORY_MODE s hodnotou CU_SHARED_MEMORY_MODE_ALLOW_OVERSIZED_SHARED_MEMORY. Komentář nad tím řádkem říká, co to stojí: vyrovnávací paměť L1 klesne na 8 kB a plánování skupin bloků musí běžet v režimu SPREAD.
Ty konstanty si Triton definuje sám, s poznámkou „remove attribute enum def once latest cuda.h is in use“. Ve veřejné hlavičce CUDA tedy zatím nejsou. Kolik sdílené paměti na multiprocesor Rubin má, se z kódu vyčíst nedá: číslo 328 kB stojí jen v popisu návrhu, v samotné změně nikde.
Na FP8 stačí polovina instrukcí
U maticového násobení Triton vychází z pravidla, že jedna instrukce zpracuje v rozměru K 256 bitů – u osmibitových operandů tedy K rovné 32. Na sm_107 se to při bloku M = 128 a K aspoň 64 zvedne na 64, u formátu MXFP4 ze 64 na 128. V testech je to vidět na poznámce, že FP8 s K = 128 vydá na sm_107 dvě instrukce MMA; na Blackwellu by jich byly čtyři. Přibylo i opakované použití operandu B: když se pod sebou počítají dva bloky M, načte se B jednou a druhá instrukce ho dostane s příznakem posledního použití.
Autor návrhu k tomu napsal „2x MMA throughput increase for FP8 and FP4“. Přepínače v kódu se ale jmenují supports2xFp8Tcgen05MMA a supports4xFp4Tcgen05MMA, a proč je u FP4 čtyřka, když se rozměr K zdvojnásobí stejně jako u FP8, z kódu nevyplývá.
LLVM o sm_107 neví
Nejzajímavější místo celé změny je nenápadné. V souboru compiler.py stojí dvakrát totéž: je-li výpočetní schopnost 107, předá se dál stovka. LLVM tedy dostane cíl sm_100 a rubinovské instrukce se do výsledku vkládají jako doslovný PTX – v testu je vidět llvm.inline_asm s řetězcem tcgen05.alloc.exclusive. Důvod k té náhradě do kódu nikdo nenapsal.
V běžné dokumentaci NVIDIE sm_107 není
Že jde o věc před vydáním, se pozná porovnáním dvou stránek téhož výrobce. Běžná dokumentace PTX končí u verze ISA 9.3 a řetězec sm_107 v ní nestojí ani jednou, kdežto sm_103 čtyřicetkrát. Vývojářský náhled pro CUDA 13.4 má navíc ISA 9.4 a ta začíná právě podporou cíle sm_107, k tomu sm_107f a sm_107a. Jméno Rubin v té dokumentaci nepadne; to je z Tritonu. Vedle sm_107 v seznamu stojí i sm_110, o kterém Triton taky ví – v TargetFeatures.h mu předepisuje zarovnání operandů FP4 stejně jako sm_100 a sm_103.
K uživatelům PyTorche to zatím nevede
Balíček torch ve verzi 2.13.0 žádá triton==3.7.1, a to na Linuxu a pro Python starší než 3.15. Dokud se ten zámek neposune, dostane novou verzi jen ten, kdo si ji doinstaluje sám. Podpora nové architektury se takhle skládá po vrstvách a každá čeká na tu pod sebou: NCCL 2.31 přidal přenos, který NVIDIA podporuje jen na Blackwellu a až od CUDA 13.3, a vLLM přeskočil PyTorch 2.12, protože přechod na dvojku třináctku ho stál čtyři ústupky v testech.
Zdroje: poznámky k vydání a návrh změny 10936 v repozitáři Tritonu, hlavička TargetFeatures.h na značce v3.8.0, dokumentace PTX ve vydané i náhledové podobě a záznamy balíčků triton a torch na PyPI.