Tekanan Register dan Spill

Register adalah penyimpanan tercepat yang dimiliki utas, tetapi setiap SM hanya memiliki jumlah yang terbatas. Seberapa banyak kernel menggunakannya disebut tekanan register.

Setiap utas yang aktif mengambil sumber daya dari satu berkas register per SM. Semakin banyak register yang diperlukan setiap utas, semakin sedikit utas yang dapat tetap aktif secara bersamaan.

Penggunaan register yang tinggi secara langsung membatasi jumlah warp yang dapat dimuat pada sebuah SM. Penurunan okupansi tersebut dapat membuat perangkat keras kekurangan pekerjaan untuk menyembunyikan latensi.

Ketika utas memerlukan lebih banyak register daripada yang tersedia, kompiler memindahkan nilai tambahan ke luar. Luapan ini disebut tumpahan register ke memori yang lebih lambat.

Nilai yang tertumpah masuk ke memori lokal, yang berada di DRAM luar cip. Setiap tumpahan menggantikan akses register bebas dengan perjalanan pulang-pergi yang lambat.

Minta kompiler melaporkan penggunaannya. Menambahkan --ptxas-options=-v ke nvcc akan mencetak jumlah register per utas dan byte tumpahan untuk setiap kernel.

Laporan tersebut mencantumkan penyimpanan dan pemuatan tumpahan. Jumlah tumpahan yang bukan nol merupakan tanda peringatan bahwa kernel Anda membayar biaya lalu lintas memori lokal yang lambat.

Anda dapat menetapkan batas saat kompilasi. Bendera --maxrregcount membatasi jumlah register per utas, dengan menukar sedikit kecepatan untuk okupansi yang lebih tinggi.

Petunjuk per kernel sering kali lebih baik. __launch_bounds__ memberi tahu kompiler ukuran blok Anda agar dapat menganggarkan register untuk okupansi yang Anda inginkan.

Sering kali Anda cukup menyimpan lebih sedikit nilai pada satu waktu. Menghitung ulang hasil yang murah atau mempersempit cakupan variabel mengurangi jumlah register yang tetap aktif.

ILP dan penggulungan perulangan meningkatkan tekanan, sedangkan pembatasan meningkatkan okupansi. Tantangannya adalah menemukan keseimbangan yang menghasilkan eksekusi tercepat, dan hanya profiler yang dapat menunjukkan letaknya.

Anda telah mempelajari bahwa tekanan register yang tinggi mengurangi okupansi dan dapat menyebabkan tumpahan ke DRAM yang lambat. Ukur penggunaan, batasi register, dan biarkan profiler memandu Anda. 🎯