Artikel ini memperkenalkan penyelidikan baru berjudul CODA: Rewriting Transformer Blocks as GEMM-Epilogue Programs, dengan matlamat utama untuk mengoptimumkan kecekapan latihan model Transformer, khususnya menyelesaikan operasi "berat memori" yang kelihatan terpisah tetapi akumulasi masa yang lama.

Penulis artikel, sumber: Machine Learning China

Pada 22 Mei, Tri Dao membagikan tweet Han Guo di media sosial. Beliau juga menulis: "Selepas beberapa penulisan semula matematik, ternyata semua komponen Transformer adalah siri GEMM + epilog (pendaraban matriks ditambah epilog). Dengan menggunakannya beberapa primitif yang dioptimakan, LLM (dan pemula) boleh menulis kernel pantas cahaya untuk semua operasi Transformer!"

Tri Dao adalah salah satu penulis utama siri FlashAttention, dan tweet ini merujuk kepada kertas kerja yang mereka terbitkan pada hari itu: CODA.

Tajuk kertas: CODA: Menulis semula Blok Transformer sebagai Program GEMM-Epilogue
Alamat kertas: https://arxiv.org/abs/2605.19269
Alamat kod: https://github.com/HanGuo97/coda-kernels

Nama ini, dibaca seperti "Finale", dilafalkan seperti "CUDA". Para penyelidik dari MIT, Princeton, Together AI, dan Meta cuba menghilangkan secara sistematik "pengiraan pecah-pecah" yang jarang diperhatikan tetapi terus menghabiskan masa dalam pelatihan Transformer, menggunakan satu abstrak pengaturcaraan baharu.

Cukai malas dalam pelatihan model besar

Untuk memahami masalah yang diselesaikan oleh CODA, terlebih dahulu perlu memahami ke mana waktu digunakan dalam pelatihan model besar.

Melatih model 1B parameter bergaya LLaMA-3 di atas satu unit NVIDIA H100, kebanyakan orang akan secara intuitif menganggap: semua masa dihabiskan untuk pendaraban matriks dan pengiraan perhatian, kerana itulah "pengiraan sebenar". Intuisi ini secara umum betul: pendaraban matriks (GEMM) dan perhatian memang mengambil bahagian terbesar dalam pengiraan.

Namun, jika anda membuka profiler dan memeriksa dengan teliti, anda akan melihat sejumlah «operator kecil» yang secara senyap menghabiskan masa: normalisasi (RMSNorm), fungsi aktivasi (SwiGLU, RoPE), penambahan sisa, pengurangan antar-lapisan... Mereka masing-masing mempunyai beban pengiraan yang kecil, tetapi sering mengangkut tensor sementara besar masuk dan keluar dari memori GPU.

Ini dikenali sebagai "bottleneck lebar pita memori": seumpama seorang chef yang sangat mahir, tetapi setiap kali ingin membuat hidangan, ia perlu membawa bahan-bahan dari gudang yang jauh, kemudian mengembalikannya selepas selesai, bukan meletakkannya di atas meja yang mudah dijangkau. Seberapa pantas tangan chef itu, masa menunggu pengangkutan tetap merupakan pembaziran yang nyata.

Lebih buruk lagi, semakin cepat pengiraan matriks menjadi dengan format presisi rendah seperti FP8 dan FP4 daripada NVIDIA, kos relatif operasi "pemindahan" ini justru meningkat: penggandaan matriks dipercepat, tetapi kos memindahkan tensor masuk dan keluar tidak berkurang secara sepadan.

Satu set data dalam kertas ini sangat jelas: Semasa melatih model dengan 1B parameter menggunakan TorchTitan di H100, operasi bukan perkalian matriks mengambil bahagian yang signifikan dalam masa perjalanan keseluruhan, dan nisbah ini akan menjadi lebih ketara apabila ketepatan FP8 diperkenalkan.

Rangka kerja pengaturcaraan yang ada hampir tidak berdaya terhadap ini. PyTorch mengungkapkan pengiraan Transformer sebagai urutan operator, dengan sempadan yang jelas antara setiap operator. Sempadan ini sangat mesra terhadap pembezaan automatik (autograd), tetapi justru menghalang pengoptimuman penggabungan antar-operator: setiap sempadan operator seringkali merupakan penulisan semula memori yang tidak perlu.

CODA: "Penutup" menyimpan harta karun

Titik tolak CODA adalah satu pemerhatian yang sederhana.

Pada GPU, satu kernel pendaraban matriks berprestasi tinggi (GEMM) dibahagikan secara struktur kepada dua bahagian: mainloop yang bertanggungjawab atas pengiraan pendaraban dan penambahan blok matriks utama, serta epilog yang melakukan penanganan akhir sebelum hasil ditulis semula ke memori video, seperti menambah bias, menukar jenis, dan penskalaan ringkas.

Makna kehadiran akhir ialah, pada masa ini, output daripada pendaraban matriks masih "hidup" dalam register pada cip, dan belum disimpan ke memori video global. Ini adalah jendela emas yang singkat: jika boleh melakukan lebih banyak pengiraan pada masa ini, ia boleh mengelakkan sepenuhnya satu perjalanan penulisan dan pembacaan semula memori video.

Wawasan utama CODA ialah: operasi yang memerlukan memori tinggi dalam Transformer sebenarnya boleh dinyatakan semula secara algebraik dan dilaksanakan dalam tetingkap "kesudahan" ini.

Ini memerlukan sedikit kecekapan matematik. Dengan contoh pola GEMM-RMSNorm-GEMM yang paling biasa: hasil darab matriks, diikuti dengan penambahan sisa, normalisasi RMS, kemudian darab matriks lain. Pendekatan tradisional melaksanakan tiga operator berasingan secara bersiri, dengan hasil sementara disimpan dua kali dalam memori GPU.

Tidak dapat diterjemahkan.

Pemparametrian semula serupa juga berlaku untuk operasi seperti SwiGLU, RoPE (Rotary Position Encoding), fungsi kerugian entropi silang, dan bahkan berlaku untuk penyebaran balik. Terdapat teorem dalam kertas ini yang membuktikan: selama bahagian hujung hadapan adalah "lokal berblok", penyebaran balik secara automatik mewarisi struktur yang sama. Sila lawati kertas asal untuk butiran lanjut.

Lima «bongkah» dan satu set «bahasa Lego»

CODA bukan satu kernel integrasi spesifik, tetapi satu set abstraksi pengaturan.

Ia mengunci litar utama GEMM yang telah dioptimakan oleh pakar, kemudian mengekspos lima jenis primitif asas yang boleh digabungkan di posisi akhir:

Pertukaran elemen demi elemen (penambahan residual, fungsi pengaktifan, RoPE)
Pemuatan dan penyimpanan vektor (pemancaran bobot RMSNorm)
Pemuatan dan penyimpanan blok matriks (menyimpan aktivasi sementara untuk penghantaran balik)
Block Reduction (Local RMS, Block Log-Sum-Exp)
Perubahan keadaan (statistik max dan sum-exp yang diperlukan untuk normalisasi dalam talian)

Dengan lima jenis blok ini, hampir semua operasi dalam penghantaran maju dan penghantaran balik Transformer standard, selain perhatian, boleh dicakup.

Yang lebih menarik ialah tingkat keluwesan sistem abstrak ini terhadap “siapa yang menulis kod”. Dalam eksperimen, kertas ini menilai dua modus pelaksanaan: satu oleh juru program manusia, dan satu lagi dihasilkan oleh Claude Code — dengan memberikan penerangan primitif CODA, beberapa contoh, dan log pelaksanaan, AI menyelesaikan sebahagian besar kod teras, dengan pengawasan ringan oleh manusia.

Kedua-dua mod menunjukkan prestasi pada tahap yang tinggi. Tri Dao dalam tweetnya mengatakan "LLM dan pengguna baru boleh menulis kernel kelajuan cahaya", yang secara langsung mencerminkan keputusan eksperimen dalam kertas ini.

Keputusan eksperimen

Pembandingan CODA memilih lawan yang lebih mencabar: cuBLAS ditambah torch.compile, serta Liger Kernel dan FlashInfer yang dioptimaskan khas untuk LLM.

Kertas ini menilai dua implementasi untuk setiap kernel: CODA (LLM) yang dihasilkan oleh Claude Code, dengan peneliti menyediakan penjelasan primitif, beberapa contoh, dan log teknik implementasi yang diperbaharui secara berterusan, di mana AI menghasilkan kod utama dan manusia memberikan pengawasan ringan; CODA (Human) yang ditulis secara bebas oleh jurutera manusia, menggunakan pendekatan reparameterisasi tingkat tinggi yang sama, tetapi tidak bergantung pada set primitif CODA itu sendiri. Hasil kedua-kedua grup dibandingkan dengan pustaka optimasi seperti cuBLAS + torch.compile, Liger Kernel, dan FlashInfer.

Pada peringkat operator tunggal, dengan contoh pola klasik GEMM-RMSNorm-GEMM, CODA mencapai prestasi yang melampaui baselins cuBLAS + PyTorch pada dimensi tersembunyi bagi tiga skala model: 1B, 7B, dan 70B. Kombinasi akhir seperti SwiGLU, RoPE, dan cross-entropy juga menunjukkan performa serupa.

Kernel yang dihasilkan oleh LLM sejajar dengan versi yang ditulis secara manual dalam sebahagian besar benchmark, dan dalam beberapa konfigurasi bahkan sedikit melebihi. Ini merupakan kesimpulan yang agak jarang berlaku dalam bidang pengoptimuman kernel GPU yang secara tradisional mempunyai rintangan yang sangat tinggi.

Keuntungan daripada penghantaran songsang sangat ketara: kernel penghantaran songsang GEMM-Residual-PartialRMS-GEMM mencapai peningkatan sebanyak 1.6 hingga 1.8 kali ganda berbanding garis dasar, manakala penghantaran songsang SwiGLU juga meningkat sebanyak kira-kira 1.4 hingga 1.6 kali ganda. Dalam arah ini, jurang antara LLM dan pelaksanaan manual juga kecil. Ini tidak mengejutkan: penghantaran songsang secara semula jadi melibatkan lebih banyak akses ke tensor sederhana, maka keuntungan daripada penggabungan akhir lebih besar; dan reka bentuk primitif CODA cukup jelas, membolehkan model AI untuk melaksanakan penggabungan dengan betul.

Dalam ujian end-to-end lapisan Transformer penuh, akselerasi forward CODA berkisar antara 5% hingga 20% pada pelbagai skala, dengan kesan yang lebih ketara pada saiz model yang lebih besar (yang sepadan dengan dimensi tersembunyi berskala 70B).

Dari segi ketepatan nombor, penyesuaian reparameterisasi CODA mengubah masa penerapan faktor penskalaan RMSNorm, tetapi eksperimen menunjukkan bahawa ralat nombornya sebanding dengan pelaksanaan rujukan PyTorch, dan dalam beberapa konfigurasi, ralatnya bahkan lebih kecil—berkat akumulator dengan ketepatan lebih tinggi dalam gelung utama GEMM.

Apa yang boleh dilakukan oleh CODA: Senarai pantas untuk menjelaskan batasan kemampuan CODA sebelum melihat perspektif yang lebih luas.

Cakupan: Hampir semua pengiraan dalam penyebaran maju dan penyebaran balik Transformer standard (seperti arsitektur LLaMA), selain perhatian dan penyematan kata, termasuk RMSNorm, penambahan sisa, aktivasi SwiGLU, kodifikasi posisi putaran RoPE, kerugian entropi silang, serta pengiraan gradien penyebaran balik untuk operasi-operasi di atas.
Kesan pecutan: Di bawah dimensi tersembunyi berskala 1B hingga 70B, peningkatan pada peringkat operator tunggal berbanding garis dasar cuBLAS + torch.compile berbeza-beza, dengan keuntungan dalam penyebaran balik paling ketara (sebahagian kernel mencapai lebih daripada 1.6 kali); pecutan end-to-end hadapan penuh pada lapisan Transformer adalah sekitar 5% hingga 20%, dengan kesan yang lebih menonjol pada model yang lebih besar.
Siapa sahaja boleh menggunakan: CODA yang diimplementasikan berdasarkan CuTeDSL (DSL Python untuk NVIDIA CUTLASS), menyokong dua cara penulisan kernel—oleh jurutera manusia dan model AI—dengan kedua-duanya mencapai prestasi tinggi.
Sekarang terhad: Saat ini hanya menyokong skenario GPU tunggal, tanpa melibatkan latihan teragih; reparameterisasi terutama ditujukan kepada arsitektur Transformer standard, kesesuaian untuk arsitektur lain masih perlu diverifikasi.

Penutup

CODA bukanlah pekerjaan yang berdiri sendiri. Ia adalah pelaksanaan konkrit daripada satu kelas pemikiran: di atas GPU, ruang pengoptimuman sebenar sering kali bukan terletak pada "apa yang dikira", tetapi pada "bagaimana ia dipindahkan".

FlashAttention membawa pengiraan perhatian ke dalam memori pada cip, sementara CODA cuba membawa fungsi normalisasi dan pengaktifan juga ke dalamnya. Triton menurunkan rintangan untuk menulis kernel tersuai, sementara ThunderKittens, TileLang, dan lain-lain meneroka ruang ini pada pelbagai peringkat. Kerja-kerja ini bersama-sama menunjuk ke arah yang sama: menyatukan kemudahan ekspresi graf operator PyTorch dengan kecekapan pelaksanaan yang hampir setara dengan CUDA tulisan tangan, dalam satu kerangka yang boleh diprogramkan.

Kalimat terakhir dalam tweet Tri Dao patut dipertimbangkan semula: “LLM dan pemula boleh menulis kernel pantas cahaya untuk semua operasi Transformer.” Di sebaliknya terdapat logik yang lebih mendalam: apabila abstrak pengaturcaraan direka dengan cukup baik, model AI itu sendiri boleh menyertai pengoptimuman infrastruktur latihan mereka. Perputaran inilah yang paling menarik dalam CODA.

Dari sudut pandang ini, nama "CODA" mungkin mempunyai makna yang lebih dalam. Dalam muzik klasik, Coda ialah bahagian penutup yang mengakhiri keseluruhan komposisi. Di sini, ia merupakan "penutup" bagi teras GEMM — dan menulis penutup ini dengan baik mungkin merupakan bab seterusnya yang penting dalam peningkatan kecekapan sistem latihan Transformer.