Tekanan Daftar dan Limpahan

Daftar ialah storan terpantas yang tersedia untuk utas, tetapi setiap SM hanya mempunyai jumlah yang terhad. Sejauh mana kernel menggunakannya dipanggil tekanan daftar.

Setiap utas yang berada dalam pelaksanaan mengambil daripada satu fail daftar bagi setiap SM. Lebih banyak daftar yang diperlukan oleh setiap utas, lebih sedikit utas yang boleh kekal serentak.

Penggunaan daftar yang tinggi mengehadkan secara langsung bilangan warp yang boleh dimuatkan pada SM. Penurunan keserentakan itu boleh menyebabkan perkakasan mempunyai terlalu sedikit kerja untuk menyembunyikan kependaman.

Apabila utas memerlukan lebih banyak daftar daripada yang tersedia, pengkompil memindahkan nilai tambahan keluar. Limpahan ini ialah limpahan daftar ke dalam memori yang lebih perlahan.

Nilai yang melimpah pergi ke memori tempatan, yang berada dalam DRAM luar cip. Setiap limpahan menggantikan capaian daftar percuma dengan perjalanan pergi balik yang perlahan.

Minta pengkompil melaporkan penggunaan. Menambahkan --ptxas-options=-v pada nvcc mencetak daftar bagi setiap utas dan sebarang bait limpahan untuk setiap kernel.

Laporan menyenaraikan stor dan muatan limpahan. Sebarang kiraan limpahan bukan sifar ialah tanda amaran bahawa kernel anda membayar kos trafik memori tempatan yang perlahan.

Anda boleh menetapkan had atas semasa kompilasi. Bendera --maxrregcount mengehadkan daftar bagi setiap utas, dengan menukar sedikit kelajuan untuk keserentakan yang lebih tinggi.

Petunjuk khusus kernel selalunya lebih baik. __launch_bounds__ memberitahu pengkompil saiz blok anda supaya ia boleh memperuntukkan daftar mengikut keserentakan yang anda inginkan.

Selalunya anda hanya perlu menyimpan lebih sedikit nilai pada satu-satu masa. Mengira semula hasil yang murah atau mengecilkan skop pemboleh ubah mengurangkan bilangan daftar yang kekal aktif.

ILP dan pengembangan gelung meningkatkan tekanan, manakala had meningkatkan keserentakan. Kemahirannya ialah mencari keseimbangan yang berjalan paling pantas, dan hanya pemprofil boleh menunjukkan kedudukannya.

Anda telah mempelajari bahawa tekanan daftar yang tinggi mengurangkan keserentakan dan boleh menyebabkan limpahan ke DRAM yang perlahan. Ukur penggunaan, hadkan daftar, dan biarkan pemprofil membimbing anda. 🎯