Hindari Konflik Bank

Memori bersama terbagi menjadi 32 bank, satu untuk setiap thread dalam warp. Menyebarkan akses ke berbagai bank memungkinkan semua 32 thread membaca dengan kecepatan penuh.

Kata berurutan berukuran 4 byte ditempatkan pada bank yang berurutan, lalu berulang kembali setelah 32. Jadi kata 0 berada di bank 0, dan kata 32 kembali berada di bank 0.

Konflik bank terjadi ketika dua thread dalam warp mengakses bank yang sama tetapi kata yang berbeda. Perangkat keras harus membuat akses tersebut berjalan secara berurutan.

Konflik N-arah mengubah satu akses cepat menjadi N akses lambat. Konflik 32-arah dapat membuat memori bersama sama lambatnya seperti jika semua akses dilakukan dalam satu antrean. 🐢

Ketika setiap thread membaca tile[threadIdx.x], setiap thread mengakses bank yang berbeda. Akses ini bebas konflik dan berjalan dengan lebar pita penuh.

Pembacaan dengan langkah 32 mengarahkan setiap thread ke bank yang sama. Ini adalah kasus terburuk, yaitu konflik 32-arah penuh.

Stride genap apa pun yang memiliki faktor persekutuan dengan 32 akan menyebabkan konflik parsial. Misalnya, stride 2 menimbulkan konflik 2-arah di seluruh warp.

Kabar baiknya, jika semua thread membaca alamat yang sama, perangkat keras menyiarkannya sekaligus. Membaca kata yang sama tidak masalah; konflik hanya terjadi jika bank-nya sama tetapi kata yang dibaca berbeda.

Untuk tile 2D, tambahkan satu kolom ekstra dengan padding +1. Ini menggeser setiap baris sehingga akses kolom tidak lagi semuanya berada di satu bank.

Kolom ekstra membuat panjang baris saling prima dengan 32. Kini, menelusuri sebuah kolom ke bawah akan mengunjungi bank yang berbeda setiap kali, sehingga konflik hilang.

Konflik bank tidak terlihat dalam kode sumber. Gunakan Nsight Compute untuk mengukur konflik memori bersama sebelum meluangkan waktu untuk memperbaikinya.

Mari kita periksa hal-hal yang membuat akses bersama berlangsung cepat.

Anda telah mempelajari bahwa memori bersama memiliki 32 bank, akses ke bank yang sama dengan kata berbeda akan diserialkan, dan padding +1 memperbaiki konflik kolom. Berikutnya: memori bersama dinamis. 🎯