Pembukaan Teknik Kernel Baru untuk GPU AMD MI300X
Redaksi akdah.ac.id – Tim MoonMath AI baru saja merilis kernel perhatian maju bf16 untuk GPU AMD MI300X. Kernel ini ditulis dalam bahasa HIP, bukan dalam assembly, dan tersedia secara open-source dengan lisensi MIT. Berdasarkan informasi yang tersedia, tim MoonMath mengklaim bahwa kernel ini mengungguli AITER v3, yaitu kernel optimasi milik AMD, dalam setiap bentuk pengujian yang dilakukan. Akses bare-metal untuk pengujian ini difasilitasi oleh HotAisle, penyedia cloud AMD.
Penjelasan tentang Kernel Perhatian
Kernels pada umumnya adalah program kecil yang dijalankan langsung di banyak inti GPU untuk melakukan perhitungan spesifik guna memaksimalkan performa hardware. Dalam konteks ini, kernel yang dirilis mengimplementasikan operasi perhatian dalam format bf16 yang hanya berjalan di MI300X. Kernel ini menerima input dalam tata letak BSHD atau BHSD tanpa perlu transposisi. Ketentuan dimensi kepala ditetapkan pada 128, dan mendukung panjang urutan apapun, termasuk perhatian silang.
Tentu ada batasan dalam implementasinya. Kernel ini tidak menyediakan masker kausal, tidak memiliki GQA, dan tidak mendukung batching varlen. Semua output dihasilkan dalam format bf16 dan berjalan eksklusif pada hardware gfx942.
Teknik Utama dan Efisiensi Memori
Teknik utama dari kernel ini bertujuan untuk menghindari dilema yang umum ditemui dalam pembuatan kode. Penggunaan intrinsic compiler memberikan kebebasan dalam penulisan kode tetapi dapat mengubah urutan atau nama operand. Sebaliknya, penggunaan assembly mentah memberikan kontrol lebih tetapi mengharuskan penanganan alamat dan register secara manual. MoonMath mengimplementasikan teknik dengan membungkus satu instruksi di dalam fungsi __device__ __forceinline__, di mana tim peneliti dapat memilih opcode. Compiler tetap akan mengalokasikan register dan melacak alur data.
Sebagian besar peningkatan performa berasal dari penempatan memori. Dalam kernel ini, data ditransfer dari HBM ke LDS, ganda-buffered, serta dibagikan oleh semua gelombang yang berjalan. Sementara itu, data untuk V disimpan dalam cache L1, dan data untuk Q serta akumulator berada di register. Pengaturan ukuran MFMA sebesar 16×16×16 dipilih bukannya 32×32×8, di mana kedua bentuk tersebut memiliki throughput yang identik. Namun, ukuran yang lebih kecil ini membantu dalam proses akumulasi, dan mengurangi beban pada akumulator, memungkinkan prefetch yang lebih dalam.
Hasil Uji dan Perbandingan dengan AITER v3
Pengujiannya dilakukan pada MI300X dengan input bf16 dan dimensi kepala 128. Pengujian ini mencakup tiga mode pembulatan: RTNE, RTNA, dan RTZ. Hasilnya menunjukkan bahwa kernel dari MoonMath mengungguli AITER v3 dari AMD dalam setiap bentuk dan mode pembulatan yang diuji. Secara geomean, kernel baru ini mencatatkan performa 1.18x lebih cepat untuk RTNE, 1.15x untuk RTNA, dan 1.08x untuk RTZ dibandingkan dengan AITER.
Khusus untuk mode RTZ yang merupakan mode tercepat AITER, bentuk (4, 16, 16384) mengalami peningkatan dari 0.95x menjadi 1.07x. Keberhasilan tersebut sangat mungkin disebabkan oleh teknik tail KV split yang diterapkan dalam kernel ini, yang diadaptasi untuk menutup kesenjangan performa.
Kesimpulan
Dengan peluncuran kernel perhatian maju bf16 ini, tim MoonMath AI menunjukkan kemampuan dan inovasi dalam pengembangan perangkat lunak untuk GPU AMD MI300X. Keunggulan dalam pengujian yang menunjukkan performa lebih baik dibandingkan kernel optimasi dari AMD memberikan harapan baru bagi para pengembang yang menggunakan platform AMD dalam proyek kecerdasan buatan dan pemrosesan data. Kernel ini tidak hanya mencerminkan kemajuan teknologi, tetapi juga berkontribusi pada komunitas open-source dengan lisensi MIT yang mendukung kolaborasi dan inovasi lebih lanjut di bidang ini.
