Kayıt Baskısı ve Taşmalar

Kayıtçılar bir iş parçacığının sahip olduğu en hızlı depolamadır, ancak her SM yalnızca sınırlı sayıda kayıtçı tutar. Bir çekirdeğin bunları ne ölçüde kullandığına kayıtçı baskısı denir.

Her yerleşik iş parçacığı, SM başına bir kayıtçı dosyasından pay alır. Her iş parçacığının ihtiyaç duyduğu kayıtçı sayısı arttıkça birlikte yerleşik kalabilen iş parçacığı sayısı azalır.

Yüksek kayıtçı kullanımı, bir SM'ye kaç warp sığacağını doğrudan sınırlar. Doluluktaki bu düşüş, gecikmeyi gizlemek için donanımda çok az iş bırakabilir.

Bir iş parçacığının mevcut olandan daha fazla kayıtçıya ihtiyacı olduğunda derleyici fazladan değerleri dışarı taşır. Bu taşma, daha yavaş belleğe gerçekleşen bir kayıtçı taşmasıdır.

Taşan değerler, yonga dışındaki DRAM'de bulunan yerel belleğe gider. Her taşma, boş bir kayıtçı erişiminin yerine yavaş bir gidiş-dönüş koyar.

Derleyiciden kullanımı raporlamasını isteyin. nvcc komutuna --ptxas-options=-v eklemek, her çekirdek için iş parçacığı başına kayıtçıları ve taşan baytları yazdırır.

Rapor, taşma depolamalarını ve yüklemelerini listeler. Sıfırdan farklı herhangi bir taşma sayısı, çekirdeğinizin yavaş yerel bellek trafiğinin bedelini ödediğine dair bir uyarıdır.

Derleme zamanında bir üst sınır belirleyebilirsiniz. --maxrregcount seçeneği, iş parçacığı başına kayıtçıları sınırlar; bunun karşılığında biraz hızdan vazgeçerek daha yüksek doluluk elde edersiniz.

Çekirdek başına verilen bir ipucu çoğu zaman daha iyidir. __launch_bounds__, derleyiciye blok boyutunuzu bildirir; böylece kayıtçıları istediğiniz doluluğa göre bütçelendirebilir.

Çoğu zaman aynı anda daha az değer tutabilirsiniz. Ucuz bir sonucu yeniden hesaplamak veya değişken kapsamını daraltmak, etkin kalan kayıtçıların sayısını azaltır.

ILP ve döngü açma baskıyı artırırken sınırlar doluluğu artırır. Asıl beceri, en hızlı çalışan dengeyi bulmaktır; bunun nerede olduğunu yalnızca profil oluşturucu söyleyebilir.

Yüksek kayıtçı baskısının doluluğu düşürdüğünü ve yavaş DRAM'e taşmalara neden olabileceğini öğrendiniz. Kullanımı ölçün, kayıtçıları sınırlayın ve profil oluşturucunun size yol göstermesine izin verin. 🎯