Tiny GPU: GPU minimal yang diimplementasikan dengan Verilog
(github.com/adam-maj)- tiny-gpu adalah implementasi GPU minimal berbasis Verilog untuk mempelajari dari dasar bagaimana GPU bekerja di level hardware, dengan fokus pada prinsip-prinsip umum GPGPU dan akselerator ML ketimbang hardware khusus grafis
- Implementasinya terdiri dari kurang dari 15 file Verilog yang terdokumentasi, dokumen arsitektur dan ISA, kernel penjumlahan serta perkalian matriks, serta dukungan simulasi kernel dan trace eksekusi
- GPU menjalankan satu kernel pada satu waktu; ia memuat program memory dan data memory, mengatur
thread_count, lalu menaikkan sinyalstartuntuk memulai kernel - Untuk penyederhanaan, tiap core memproses satu block pada satu waktu, dan tiap thread memiliki ALU, LSU, PC, serta register file sendiri, tetapi diasumsikan semua thread berkumpul kembali ke PC yang sama setelah setiap instruksi
- Sebagian besar fitur GPU modern seperti cache bertingkat, shared memory, memory coalescing, pipelining, warp scheduling, branch divergence, dan barrier dikeluarkan demi memprioritaskan struktur untuk pembelajaran
Masalah yang ingin dipecahkan tiny-gpu
- Ada banyak materi untuk mempelajari CPU, mulai dari arsitektur hingga sinyal kontrol, tetapi detail teknis level rendah GPU modern sebagian besar tetap proprietary karena pasar yang kompetitif
- Ada banyak materi pemrograman GPU, tetapi hampir tidak ada materi untuk mempelajari bagaimana GPU bekerja di level hardware
- Implementasi GPU open source seperti Miaow dan VeriGPU menargetkan kelengkapan fitur dan operasi, sehingga strukturnya kompleks
- tiny-gpu menghilangkan banyak kompleksitas kartu grafis kelas produksi dan berfokus pada elemen inti yang umum pada akselerator hardware modern
- Komponen penting dalam arsitektur GPU
- Cara model pemrograman SIMD diimplementasikan pada hardware
- Cara GPU menangani bandwidth memori yang terbatas
Arsitektur keseluruhan
- tiny-gpu dirancang untuk menjalankan hanya satu kernel pada satu waktu
- Prosedur eksekusi kernel adalah sebagai berikut
- Memuat kode kernel ke global program memory
- Memuat data yang diperlukan ke data memory
- Menentukan jumlah thread yang akan dijalankan pada device control register
- Mengatur sinyal
startke high untuk menjalankan kernel
- GPU terdiri dari unit-unit berikut
- Device control register
- Dispatcher
- Jumlah compute core yang dapat bervariasi
- Memory controller untuk data memory dan program memory
- Cache
Eksekusi kernel dan pembagian thread
- Device control register berperan menyimpan metadata eksekusi kernel; pada tiny-gpu, yang disimpan hanya
thread_count, yaitu jumlah total thread yang akan dijalankan - Dispatcher membagikan thread ke beberapa compute core saat kernel dimulai
- Membentuk kumpulan thread yang dapat dijalankan paralel sebagai block
- Mengirim block ke core yang tersedia untuk diproses
- Memberi tahu bahwa eksekusi kernel selesai setelah semua block diproses
- Core yang disederhanakan memproses satu block pada satu waktu
- Tiap thread memiliki ALU, LSU, PC, dan register file khusus
- Mengelola eksekusi instruksi thread di atas resource ini adalah salah satu masalah sulit pada GPU
Struktur memori dan controller
- GPU dibuat untuk berinteraksi dengan external global memory, dan untuk penyederhanaan, data memory dan program memory dipisahkan
- Spesifikasi data memory
- Pengalamatan 8-bit
- Total 256 baris
- Data 8-bit
- Tiap baris menyimpan nilai kurang dari 256
- Spesifikasi program memory
- Pengalamatan 8-bit
- Total 256 baris
- Data 16-bit
- Sesuai ISA, tiap instruksi berukuran 16-bit
- Memory controller melacak request memori dari core, membatasi request sesuai bandwidth external memory sebenarnya, dan mengirimkan respons ke resource yang benar
- Tiap memory controller memiliki jumlah channel tetap sesuai bandwidth global memory
- Cache adalah fitur yang sedang dikerjakan; cache menyimpan data yang diambil dari external memory ke SRAM di dalam device agar request berikutnya dapat mengambilnya lebih cepat dan bandwidth memori dapat dipakai untuk data baru
Komposisi internal core
- Tiap core memiliki satu scheduler yang mengelola eksekusi thread
- Scheduler tiny-gpu mengeksekusi instruksi satu block hingga selesai lalu mengambil block baru, dan menjalankan instruksi semua thread dalam urutan yang tersinkronisasi
- Pada scheduler yang lebih maju, pipelining dan warp scheduling dapat meningkatkan utilisasi resource
- Kendala utama scheduler muncul dari latensi saat memuat dan menyimpan data dari global memory
- Sebagian besar instruksi dapat dijalankan secara sinkron
- Operasi load-store seperti
LDRdanSTRbersifat asinkron, sehingga eksekusi instruksi harus disusun dengan mempertimbangkan waktu tunggu yang panjang
- Fetcher mengambil instruksi pada program counter saat ini dari program memory secara asinkron
- Decoder mendekode instruksi yang diambil menjadi sinyal kontrol untuk eksekusi thread
- Register file tiap thread menyimpan data yang sedang dihitung dan memungkinkan pola SIMD
- Register read-only berisi
%blockIdx,%blockDim,%threadIdx - Kernel dapat dijalankan dengan data berbeda sesuai ID thread lokal
- Register read-only berisi
- ALU tiap thread memproses instruksi aritmetika
ADD,SUB,MUL,DIV CMPmengeluarkan apakah hasil selisih dua register bernilai negatif, nol, atau positif, dan menyimpan hasilnya di registerNZPmilik unit PC- LSU tiap thread mengakses global data memory serta menangani
LDR,STR, dan latensi memori asinkron - PC tiap thread menentukan instruksi berikutnya yang akan dieksekusi
- Secara default meningkat 1 pada setiap instruksi
BRnzpbercabang ke baris program memory tertentu jika kondisi registerNZPyang ditetapkan olehCMPsebelumnya terpenuhi- Loop dan conditional diimplementasikan dengan cara ini
- Untuk penyederhanaan, tiny-gpu mengasumsikan semua thread berkumpul kembali ke PC yang sama setelah setiap instruksi
- Pada GPU nyata, thread individual dapat bercabang ke PC yang berbeda, dan saat itu kelompok thread yang sebelumnya diproses bersama terpecah menjadi beberapa alur eksekusi; ini disebut branch divergence
ISA
- tiny-gpu mengimplementasikan ISA 11 instruksi untuk menjalankan kernel sederhana proof-of-concept seperti penjumlahan matriks dan perkalian matriks
- Instruksi yang didukung
BRnzp: lompat ke baris program memory lain jika kondisiNZPterpenuhiCMP: membandingkan dua nilai register dan menyimpan hasilnya ke registerNZPADD,SUB,MUL,DIV: operasi aritmetika dasar untuk matematika tensorLDR: memuat data dari global memorySTR: menyimpan data ke global memoryCONST: memuat nilai konstanta ke registerRET: sinyal penghentian eksekusi thread saat ini
- Tiap register ditentukan dengan 4-bit, sehingga total ada 16 register
- 13 register dari
R0hinggaR12adalah register bebas yang dapat dibaca dan ditulis - 3 terakhir adalah register khusus read-only yang menyediakan
%blockIdx,%blockDim,%threadIdxyang dibutuhkan SIMD
- 13 register dari
Alur eksekusi
- Saat mengeksekusi instruksi, tiap core mengikuti alur kontrol tahap berikut
FETCH: mengambil instruksi berikutnya dari PC saat iniDECODE: mendekode instruksi menjadi sinyal kontrolREQUEST: meminta data ke global memory jikaLDRatauSTRmembutuhkannyaWAIT: menunggu respons global memory jika diperlukanEXECUTE: menjalankan kalkulasi atas dataUPDATE: memperbarui register file dan registerNZP
- Alur kontrol ini disusun demi kesederhanaan dan kemudahan dipahami
- Dalam implementasi nyata, sebagian tahap dapat dipadatkan untuk mengoptimalkan waktu pemrosesan, atau pipelining dapat mengoordinasikan eksekusi beberapa instruksi pada resource core
- Tiap thread melakukan kalkulasi pada data dalam register file khususnya dengan mengikuti jalur eksekusi yang sama
- Mirip dengan diagram CPU, tetapi berbeda karena
%blockIdx,%blockDim,%threadIdxberada di register read-only sehingga memungkinkan fungsi SIMD
Contoh kernel
- Untuk proof-of-concept ISA, kernel penjumlahan matriks dan perkalian matriks ditulis
- File test di repository dapat mensimulasikan kernel-kernel ini sepenuhnya pada GPU serta menghasilkan status data memory dan trace eksekusi penuh
-
Penjumlahan matriks
matadd.asmmenjumlahkan dua matriks 1 x 8- Penjumlahan per elemen pada 8 elemen dilakukan masing-masing di thread terpisah
- Menunjukkan pemrograman SIMD menggunakan register
%blockIdx,%blockDim,%threadIdx - Mencakup manajemen memori asinkron menggunakan instruksi
LDRdanSTR
-
Perkalian matriks
matmul.asmmengalikan dua matriks 2 x 2- Menghitung dot product baris dan kolom terkait untuk tiap elemen
- Menunjukkan percabangan di dalam thread menggunakan
CMPdanBRnzp - Semua cabang berkumpul kembali, sehingga bekerja pada implementasi tiny-gpu saat ini
Simulasi
- Untuk menjalankan simulasi kernel, diperlukan iverilog dan cocotb
- Langkah persiapan
- Instal compiler Verilog dan cocotb dengan
brew install icarus-verilogdanpip3 install cocotb - Unduh versi terbaru sv2v, ekstrak, lalu tambahkan binary ke
$PATH - Jalankan
mkdir builddari root repository
- Instal compiler Verilog dan cocotb dengan
- Simulasi kernel dijalankan dengan
make test_matadddanmake test_matmul - Hasil eksekusi ditulis sebagai file log di
test/logs- Status awal data memory
- Trace eksekusi penuh kernel
- Status akhir data memory
- Di awal tiap file log terlihat matriks input, dan pada data memory akhir di bagian akhir terlihat matriks hasil
- Trace eksekusi mencakup status eksekusi semua thread di semua core pada tiap cycle
- Instruksi saat ini
- PC
- Nilai register
- Informasi status
Fitur GPU lanjutan yang sengaja tidak disertakan
- Demi penyederhanaan, tiny-gpu mengecualikan sebagian besar elemen peningkatan performa dan fungsi GPU modern
-
Cache bertingkat dan shared memory
- GPU modern menggunakan beberapa layer cache untuk mengurangi akses global memory
- tiny-gpu hanya mengimplementasikan satu layer cache yang menyimpan data terbaru antara resource peminta dan memory controller
- Cache multi-layer mengurangi waktu load dengan menyimpan data yang sering dipakai lebih dekat ke lokasi pemakaiannya
- GPU juga dapat menggunakan shared memory agar thread dalam block yang sama dapat saling bertukar hasil bersama
-
Memory coalescing
- Beberapa thread yang berjalan paralel sering mengakses alamat berurutan, seperti elemen matriks yang berdekatan
- Memory coalescing menganalisis request memori yang menumpuk di queue dan menggabungkan request yang berdekatan menjadi satu transaksi
- Tujuannya mengurangi waktu yang dipakai untuk pengalamatan dan memproses request bersama-sama
-
Pipelining
- Core tiny-gpu baru memulai instruksi berikutnya setelah eksekusi satu instruksi dari satu kumpulan thread selesai
- GPU modern men-stream eksekusi beberapa instruksi berurutan sambil tetap menjamin instruksi yang memiliki dependency dijalankan secara berurutan
- Ini meningkatkan utilisasi resource agar resource core tidak menganggur dalam situasi seperti menunggu request memori asinkron
-
Warp scheduling
- Block dibagi menjadi warp, yaitu batch thread yang dapat dijalankan bersama
- Saat satu warp sedang menunggu, instruksi warp lain dijalankan sehingga beberapa warp dapat diproses bersamaan pada satu core
- Mirip dengan pipelining, tetapi menangani instruksi dari thread yang berbeda
-
Branch divergence
- tiny-gpu mengasumsikan semua thread dalam satu batch berada pada PC yang sama setelah tiap instruksi
- Pada kenyataannya, thread individual dapat bercabang ke baris berbeda tergantung data
- Thread dengan PC berbeda dipisahkan menjadi alur eksekusi tersendiri, dan titik saat mereka berkumpul kembali juga harus dikelola
-
Sinkronisasi dan barrier
- GPU modern dapat menetapkan barrier agar grup thread dalam block yang sama menunggu hingga semuanya mencapai titik tertentu
- Ini berguna saat thread perlu bertukar data bersama, untuk menjamin pemrosesan data sudah selesai
Pekerjaan berikutnya
- Item peningkatan ke depan adalah sebagai berikut
- Menambahkan instruction cache sederhana
- Membangun adapter agar GPU dapat digunakan di Tiny Tapeout 7
- Menambahkan branch divergence dasar
- Menambahkan memory coalescing dasar
- Menambahkan pipelining dasar
- Mengoptimalkan alur kontrol dan penggunaan register untuk meningkatkan cycle time
- Menulis kernel grafis dasar atau menambahkan hardware grafis sederhana untuk menunjukkan fungsi grafis
- Pengguna yang ingin memperbaiki repository dapat berkontribusi lewat PR
1 komentar
Komentar Hacker News
Sebagai pengecualian, Intel menerbitkan banyak dokumentasi teknis GPU: https://kiwitree.net/~lina/intel-gfx-docs/prm/
Manual i810/815 juga bisa ditemukan online, dan selain celah aneh pada periode sebelum 965—dengan 855/910/915/945 yang hilang—dokumentasinya terbilang cukup konsisten
Ini bahkan mencakup dokumen arsitektur set instruksi untuk produk saat ini dan lama, tetapi tampaknya lebih ditujukan bagi implementor daripada penjelasan tingkat tinggi untuk penggemar yang tertarik
Semoga semua perusahaan mengikuti pendekatan ini
Namun menurut saya ini lebih dekat ke koprosesor SIMD
Untuk disebut GPU, menurut saya setidaknya harus ada semacam output display
Saya tahu istilahnya sekarang cukup longgar karena Nvidia dan lainnya belakangan juga menjual varian arsitektur grafis khusus server sebagai GPU, tetapi dalam desain GPU, bagian grafis masih menyumbang porsi kompleksitas yang besar
GPU yang tidak menampilkan output pun tetap berguna
Di tempat kerja saya ada sekitar 75 workstation dengan Quadro kelas menengah, tetapi kartunya hanya punya mini-DisplayPort sementara perusahaan hanya membelikan kabel HDMI, jadi semuanya terhubung ke grafis terintegrasi
Meski begitu, kartu-kartu itu tetap mengakselerasi software dan memproses grafis; hanya saja tidak mengeluarkan tampilan ke layar
Ada contoh lain juga: https://github.com/jbush001/NyuziProcessor
Kira-kira perlu volume sebanyak apa agar TSMC atau foundry lain bisa memproduksi prosesor semacam ini secara ekonomis?
Saya ingin mencoba FPGA, tetapi jujur saja bahkan sulit membayangkan harus mulai dari mana, dan seluruh bidangnya terasa cukup mengintimidasi
Target akhir saya adalah membuat kartu akselerator untuk LLM, dan meski itu tujuan yang sepenuhnya saya tetapkan secara acak, sepertinya akan banyak tumpang tindih dengan proyek ini; mungkin bedanya hanya pada bagian memory offloading untuk memuat model yang lebih besar
Memulai FPGA harus dipecah menjadi beberapa subketerampilan, dan ekspektasinya juga perlu disesuaikan
Kita tidak mengharapkan seorang software engineer sejak awal membuat seluruh komputer dari prinsip dasar, menulis arsitektur set instruksi, memahami bahasa mesin, mengubahnya menjadi assembly, lalu bahkan mengembangkan bahasa pemrograman untuk membuat aplikasi dengan kode Python
Yang benar adalah mulai dari atas lalu turun ke bawah dalam stack
Jika fokus pada mengabstraksikan kompleksitas dan membangun sistem dengan IP yang sudah jadi, desain FPGA cukup mudah
Biasanya saya merekomendasikan sesuatu seperti MATLAB, karena Anda bisa membuat aplikasi awal dengan HDL Coder pada DevKit yang memiliki reference design
Kalau tidak, bebannya sangat besar: mempelajari arsitektur komputasi digital, Verilog, timing, transceiver/I/O, perencanaan pin, Quartus/Vivado, simulasi/verifikasi, sistem embedded, dan sebagainya
Singkatnya, mulailah dari desain tingkat sistem, pelajari cara mengambil IP plug-and-play dan menghubungkannya di level teratas, lalu masukkan modul itu ke dalam desain referensi yang sudah dibuat
Setelah itu, perlahan kupas lapisan-lapisannya untuk menyingkap kompleksitas di bawahnya
Anda bisa mengerjakan latihan sambil membaca Harris&Harris
Setelah menyelesaikan latihan dan bukunya, saatnya mulai proyek sendiri
Sebagai catatan, HackerMojo juga mengadakan pertemuan mingguan, dan meski tidak berada di Valley, Anda bisa ikut lewat Zoom
Mulailah dari sana
Hardware sebenarnya bisa dibilang bagian yang mudah, sekaligus bagian yang sulit dari sisi manufaktur
alwayssekuensial di sini?Ada situs bernama opencores yang mengumpulkan berbagai proyek HDL open-source
Penasaran apakah sekarang ada simulator HDL terdistribusi skala besar setingkat HPC yang cukup bagus
Memanfaatkan GPU modern untuk simulasi tingkat RTL tampaknya masuk akal
Apakah itu bukan situs yang sama, melainkan tempat lain yang mirip?
Apakah core CUDA modern biasanya punya pembagian sebagai instruksi sungguhan, atau biasanya diemulasikan dengan software?
Rangkaian pembagian hardware nyata memakan area sangat besar, jadi saya tidak menyangka itu akan masuk ke ALU GPU
Di Verilog memang sangat mudah menulis satu baris
DIV: begin alu_out_reg <= rs / rt; end, tetapi satu baris itu menghabiskan banyak silikonKalau hanya mensimulasikan Verilog, fakta itu mungkin tidak terlihat
Proyeknya berhenti di simulasi, dan untuk menjadikannya hardware nyata perlu jauh lebih banyak pekerjaan
Secara pribadi, menurut saya hal seperti ini sebaiknya disebut dengan nama lain
Jarak di antara keduanya makin menyempit dan keduanya saling menambahkan fungsi satu sama lain, tetapi tetap ada perbedaan yang cukup besar
Menurut saya ini berkaitan dengan Hukum Amdahl [0]
Dalam arti itu, CPU bisa disebut prosesor yang dioptimalkan untuk latensi, sedangkan GPU prosesor yang dioptimalkan untuk throughput
Lebih spesifik lagi, [1] CPU juga bisa disebut prosesor dengan dependensi data yang panjang dan dalam, sementara GPU prosesor dengan dependensi data yang lebar dan datar
[0]: https://en.wikipedia.org/wiki/Amdahl%27s_law
[1]: https://en.wikipedia.org/wiki/Data_dependency
Tensor hanyalah array berdimensi n
Di atasnya bisa dipasang software atau firmware agar berperilaku seperti GPU
Saya mencoba menyusun potongan-potongannya dari source EDK2, tetapi tidak jelas seberapa banyak yang spesifik untuk QEMU
tiny-gpubahwa semua thread “berkonvergensi” ke program counter yang sama setelah setiap instruksi adalah penyederhanaan yang terlalu naifPada GPU nyata, thread individual bisa bercabang ke PC yang berbeda, dan muncul divergensi cabang ketika grup thread yang awalnya diproses bersama terpecah menjadi eksekusi terpisah
Mungkin sebaiknya mencoba pemrograman GPU dulu sebelum membuat GPU di silikon
Selain itu, menyebutnya SIMD pun rasanya kurang tepat
Orang ini adalah orang yang dulu menyambung-nyambungkan rangkaian orang lain untuk membuat LED berkedip lalu mengatakan telah membuat CPU
__syncthreads()di setiap eksekusi?