Kayıt ve Paylaşılan Bellek Sınırları

Her SM'nin sabit bir kayıtçı ve paylaşılan bellek havuzu vardır. Her yerleşik blok, kendisine düşen payı bu havuzlardan almalıdır.

Kernel'iniz iş parçacığı başına belirli sayıda kayıtçı kullanır. Bunu blok başına iş parçacığı sayısıyla çarparsanız, bir bloğun kayıtçı maliyetini görürsünüz.

Her iş parçacığı çok sayıda kayıtçıya ihtiyaç duyarsa daha az iş parçacığı sığar ve daha az blok yerleşik kalır. Kayıtçı kullanımı yüksek kernel'ler doluluk kaybeder.

Bu bayrakla derleme yaptığınızda nvcc, iş parçacığı başına kayıtçı sayısını yazdırır ve bir kernel'in kayıtçıya fazla yüklenip yüklenmediğini görmenizi sağlar.

Bir başlatma sınırıyla üst sınır belirleyebilir; böylece derleyiciyi daha az kayıtçı kullanmaya ve daha fazla warp'ı yerleşik tutmaya yönlendirebilirsiniz.

Her blok paylaşılan bellek isteyebilir. SM, paylaşılan bellek havuzunun izin verdiği kadar bloğu sığdırır; bu miktar çoğu zaman 48 ila 100 KB arasındadır.

Büyük bir __shared__ döşemesi isterseniz SM başına yalnızca bir veya iki blok sığar. Büyük döşemeler, veri yeniden kullanımını artırmak karşılığında doluluğu azaltır.

SM; kayıtçılara, paylaşılan belleğe ve warp sınırına göre izin verilen blok sayısını hesaplar. Bu değerlerin en küçüğü doluluğu belirler.

Bir iş parçacığı izin verilen sayıdan daha fazla kayıtçıya ihtiyaç duyduğunda fazlalıklar yavaş yerel belleğe taşar. Taşmalar, düşük doluluktan daha fazla zarar verebilir.

Kayıtçıları veya paylaşılan belleği azaltmak doluluğu artırır, ancak aşırı agresif bir azaltma taşmalara yol açar. En uygun nokta bir dengedir.

Ayarlama yapmadan önce gerçek kayıtçı ve paylaşılan bellek kullanımını her zaman derleyiciden okuyun. Tahmin etmek genellikle yanlış ayarı seçmenize neden olur.

Kayıtçıların ve paylaşılan belleğin sabit SM bütçeleri olduğunu, en sıkı sınırın da doluluğu belirlediğini gördünüz. Taşmalara dikkat edin. 🧮